Clarification needed for structs with generic pointers as kernel arguments.

[karolherbst]

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.

[pierremoreau]

I agree that the specification does not mention any particular restrictions on pointers inside structures, that are passed as argument to kernels. In other parts of the specification, for example in Section 6.9 of the OpenCL C 2.0 Specification:

Arguments to kernel functions in a program cannot be declared with the built-in scalar types bool, size_t, ptrdiff_t, intptr_t, and uintptr_t or a struct and/or union that contain fields declared to be one of these built-in scalar types.

I think though you could have in practice pointer to the constant address space as well, or even to local (for example, if instead of passing a pointer to local and pointer to constant as separate arguments, you just group them in a structure and pass that to the kernel). So IMO, the only thing needed would be to say in the specification that the restrictions on pointers passed as argument to kernels, also applies to pointers inside structures passed as argument to kernels.

[karolherbst]

@pierremoreau but this makes parsing the the input data more challanging, because how to you differ between a SVM pointer inside a struct and a non SVM memory handle?

[pierremoreau]

Non-SVM handles would be decorated with constant, local or global, whereas SVM handles would be undecorated; I am not sure whether you are allowed to have local SVM handles and global SVM handles, but that would be really weird IMHO. These are my thoughts, not necessarily what the spec says.

[karolherbst]

no. SVM doesn't change anything in the kernel. You just have a shared address space in the application and on the GPU. I think it is even defined that SVM only covers global memory.

[pierremoreau]

Okay, good, that keeps things relatively simple. :smiley: