OpenCL kernel fails to compile asking for address space qualifier

Question

The following opencl code fails to compile.

typedef struct {

    double d;
    double* da;
    long* la;
    uint ui;

} MyStruct;


__kernel void MyKernel (__global MyStruct*  s) {

}

The error message is as follows.

line 11: error: kernel pointer arguments must point to addrSpace global, local, or constant
  __kernel void MyKernel (__global MyStruct*  s) {
                                              ^

As you can see I have clearly qualified the argument with '__global' as the error suggests I should. What am I doing wrong and how can I resolve this error?

Obviously this happens during kernel compilation so I haven't posted my host code here as it doesn't even get further than this.

Thanks.

__kernel void MyKernel (__global struct MyStruct* s) should it not have the struct keyword in between, seeing this is C? — Tony The Lion
@TonyTheLion I don't think so, it's been typedef'ed as MyStruct, so the struct is implicit. I think the problem is that he has pointers in his struct, which is not allowed. Why do you have pointers in your struct, what is their meaning? You cannot point to host memory from your kernel like that. — Thomas
@TonyTheLion Adding the struct keyword does actually make the kernel compile successfully. But I'm not sure if that's doing what I want it to do. I'll explain in my next comment to Thomas. — junkie
@Thomas The above struct is intended to have a double, a double array, a long array and an unsigned int. In my host code I have an array of three of these structs (though could be more) that I'm trying to pass to the kernel as an argument and access each one using array indices. As a side note, adding the struct keyword does actually make the kernel compile but semantically is that doing what I want it to do? — junkie
@junkie Hmm, maybe Tony was right then, though if my C-fu is correct, struct should not be needed. Does the kernel actually work now? Or just compile? I've never tried having arrays inside structs, I do not see how it would work since a pointer is mapped to a cl_mem host-side, and your struct is already a cl_mem, but I could be wrong. — Thomas

Thomas Thomas · Accepted Answer · 2012-09-25T00:13:30

I think the problem is that you have pointers in your struct, which is not allowed. You cannot point to host memory from your kernel like that, so pointers in kernel argument structs don't make much sense. Variable-sized arrays are backed up in OpenCL by a cl_mem host object, and that counts for one whole argument, so as far as I know, you can only pass variable-sized arrays directly as a kernel argument (and adjust the number of work units accordingly, of course).

You might prefer to put size information in your struct and pull out the arrays as standalone kernel arguments.

OpenCL kernel fails to compile asking for address space qualifier

1 Answers