OpenCL struct argument invalid address space

Go To


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;                                               
2012-04-04 19:56
by gkaran89


global is a pointer qualifier (address space), so you have to pass global const t_param* params.

2012-04-04 20:32
by eudoxos
Thanks. That was indeed the problem. I was testing it before seeing your reply and came back to comment that I found the solution - gkaran89 2012-04-04 20:35
I found it nice in general (no idea about this case) to see errors from different compilers (AMD, Intel, for instance); not all of them are helpful in all cases - eudoxos 2012-04-04 20:37