OpenCL struct argument invalid address space

Question

I'm a total beginner with OpenCL and I'm trying to make the following kernel to work. When I run the program I get an error in the build process of the kernel program. More specificity the error is the following:

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,
                                                               ^

You can see the kernel here. In the beginning I though that it was because I hadn't the structs defined inside the kernel but even when I added them the problem still exists. What am I doing wrong here?

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 is a pointer qualifier (address space), so you have to pass global const t_param* params.