How to overcome Error Code -11 cl_build_program_failure

Question

I am trying to outsource functions on linked lists (simulation of particles) to OpenCL kernels, using shared virtual memory. I tried starting by simply iterating through a linked list and changing one value of every element in it (structs).

This is the .cl file (the typedef to real is to keep consistent with host code):

//real is of type cl_double
typedef cl_double real;
typedef cl_double2 real2;

typedef struct
{

//  Mass
real m;
//  Position
real2 x;
//  Velocity
real2 v;
//  Force
real2 F;
//  Force_old
real2 F_old;
//  Bodytype
cl_char body;

} Particle;

//  Datastructure of linked list
typedef struct ParticleList
{

Particle p;
struct ParticleList *next;
} ParticleList;

//  A cell is defined as the anchor of a linked list
typedef ParticleList *Cell;

__kernel void test(
__global ParticleList *pList){

 //  Check if pList->next is NULL
if(pList->next != NULL){

    while(pList->next != NULL){

        pList->p.body = 'Z';
        pList = pList->next;
   }
}


}

Any idea on why it does not compile the .cl file? As I've understood I can define structs, typedefs and functions in the source code and use them in a kernel function.

clCreateProgramWithSource returns CL_SUCCESS, but clBuildProgram on that program returns error code -11.

Maybe some tips for debugging opencl c?

EDIT: calling clGetProgramBuildInfo yields:

1:49:19: error: assigning 'struct ParticleList *__global' to '__global 
ParticleList *' (aka '__global struct ParticleList *') changes address space
  of pointer
        pList = pList->next;
              ^ ~~~~~~~~~~~

I'm not sure what that means, can I not dereference pointers that are in the devices adress space?

Your first step is to call clGetProgramBuildInfo. This should give you specific source code lines and error messages for tracking down your problem. (I'm unfortunately not familiar with OpenCL 2.0/SVM, so I can't say for sure what the issue is just by looking at your code.) Add the output from clGetProgramBuildInfo to your question in if any doubt about what it means. — pmdj
Thank you, that makes debugging a lot easier. I updated the post ith the compiler error — M408

pmdj pmdj · Accepted Answer · 2018-11-17T10:21:53

Pointers always refer to a specific address space: global, constant, local, or private. Even if the pointer is not annotated, one of those is picked by default depending on context. In your case,

__global ParticleList *pList

Is (correctly) annotated as being in the global space, while the field next in your struct has no annotation:

struct ParticleList
{
  Particle p;
  struct ParticleList *next; // <--- no address space specified, defaults to `private`
}

Obviously, the next field does not point to a struct allocated in private memory, so this default is incorrect, you should explicitly specify global.

(Personally, I feel that defaulting to an address space was a mistake in the design of OpenCL, and it should always be explicit, but what can you do.)

How to overcome Error Code -11 cl_build_program_failure

1 Answers