kernel argument without memory space specifier


I’d like to ask, where a kernel argument is stored, when it doesn’t have _global, _public or other memory specifier?

Here is a sample code from one presentation:

char* source = "__kernel void MyKernel(__global int* buffer, int N) {
               "    if (get_global_id(0) < N) buffer[get_global_id(0)] = 7;

Parameter N does not have memory space specifier. Where is it stored?

I believe the default is private.

That makes sense…

but how does GPU handle private memory objects? Are they copied to GPU with every thread instance, resulting in extensive brandwidth usage? (f.e. in this case, provided we have a large matrix) Would’t it be better to tag N as “_constant”?

I don’t think that is defined in the spec. I would play it safe and define such things as constants to avoid the ambiguity.