Clarification needed for structs with generic pointers as kernel arguments.

Starting with OpenCL 2.0 one can declare generic pointers by leaving out the address space annotation. The current specification disallows having generic pointers or arrays declared as kernel arguments.

the OpenCL C specification states (6.5 in 2.0, 1.5 in 2.2): "kernel function arguments declared to be a pointer or an array of a type must point to one of the named address spaces global, local or __constant."

But having a struct with generic pointers are perfectly valid for kernel parameters:

struct NumberList {
        struct NumberList *next;
        unsigned int val;
};

__kernel void test(__global struct NumberList *list)
{
        struct NumberList temp;
        temp.next = 0;
        temp.val = 50;
        list->next = &temp;
}

This can come up for example if a kernel uses the fine grained system SVM features of OpenCL and wishes to share linked list between the host and the kernel.

In practice the pointer can only have a value stored to the global address space, but maybe it should be forbidden in the first place. Anyway, I think the specification should clear up this situation and have it well defined.

KhronosGroup / OpenCL-Registry

Clarification needed for structs with generic pointers as kernel arguments. #47