OpenCL struct argumento espacio de direcciones no válido

Question

Soy un principiante total con OpenCL y estoy tratando de hacer que el kernel siguiente funcione. Cuando ejecuto el programa, aparece un error en el proceso de compilación del programa kernel. Más específico el error es el siguiente:

Error: Failed to build program executable! 
<program source>:19:64: error: invalid address space for argument to __kernel function 
__kernel void accelarate_flow(__global const   t_param params, 
                  ^

Puede ver el kernel aquí. Al principio pensé que era porque no tenía las estructuras definidas dentro del kernel, pero incluso cuando las agregué, el problema todavía existe. ¿Qué estoy haciendo mal aquí?

const char *accelerate_flow_kernel_source = 
#pragma OPENCL EXTENSION cl_khr_fp64 : enable         
typedef struct                 
{                    
    int nx;                 
    int ny;                 
    int maxIters;                
    int reynolds_dim;               
    double density;                
    double accel;                
    double omega;                
} t_param;                  

typedef struct                 
{                    
    double speeds[9];               
} t_speed;                  

__kernel void accelarate_flow(__global const   t_param params,   
           __global const   int*  obstacles,  
           __global    t_speed* cells,   
             const unsigned int  count)   
{                    
    int pos = get_global_id(0);             
    if(pos >= count || pos % params.nx != 0) return;        
    double w1,w2;                
    w1 = params.density * params.accel/9.0;         
    w2 = params.density * params.accel/36.0;         
    if(!obstacles[pos] &&              
    (cells[pos].speeds[3] - w1) > 0.0 &&          
    (cells[pos].speeds[6] - w2) > 0.0 &&          
    (cells[pos].speeds[7] - w2) > 0.0)          
    {                   
    cells[pos].speeds[1] += w1;            
    cells[pos].speeds[5] += w2;            
    cells[pos].speeds[8] += w2;            
    cells[pos].speeds[3] -= w1;            
    cells[pos].speeds[6] -= w2;            
    cells[pos].speeds[7] -= w2;            
    }                   
}                    

Answer 1

global es un calificador de puntero (dirección de espacio), por lo que tiene que pasar global const t_param* params.