OpenCL struct argument invalid address space
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;                                               
      }                                                                           
    }                                                                             

OpenCL struct argument invalid address space
===

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;                                               
      }                                                                           
    }                                                                             
