How many copies of a global variable declared inside an opencl kernel function is maintained in the global address space

Question

I'm new to Opencl programming. To learn opencl better, after spending some time reading some tutorials, I started developing a simple pattern matching kernel function. But I have some doubts:

First, I have global variables declared inside the kernel function. Does it mean every work item shares a single copy of each variable?

Second, how can I use the standard C libraries, esp. "string.h".

   __kernel void matchPatterns_V1(__global char *strings, __global char *patterns, __global int *matchCount,
                            int strCount, int strLength, int patCount, int patLength) {


    int id = get_global_id(0);
    int rowIndex = id*strLength;
    int i, matches = 0;     

    __global char *pos = strings;
    __global char *temp = strings;
    __global char *pat = patterns;

    for(i = 0; i < patCount; i++)
    {
            temp = &strings[rowIndex];      
            pat = &patterns[i*patLength];
            while(pos != '\0') {
                    pos = StrStr(temp, pat);
                    if(pos != '\0') {
                            matches++;
                            temp = pos + patLength;
                    }
            }
    }
    matchCount[id] = matches;
    }

To summarize, does each work item has its own copy of the variables 'pos', 'temp', and 'pat'?

Any advice in learning Opencl is highly appreciated, including suggestion for best books/tutorial sites.

Answer 1

No, it's in global memory space so there is, in general, a single copy per kernel invocation, to be shared by all work items. Writing to global memory is dangerous if you cannot guarantee each work group has its own unique "item" in global memory - or, more generally, that no two work items in your kernel will write to the same location in memory at the same time) because there will be race conditions.

If you are simply reading data from those global memory variables, it doesn't matter, of course.

You also cannot declare __global variables inside your kernel as they can only be kernel arguments. If you attempt to do so, you will get the following from the OpenCL compiler:

error: global variables cannot be allocated inside kernel code

For good reason, barring technical impossibilities: a global variable would serve no purpose at all... the only possible reason I can even think of would be communicating between work items, which would be an insane design pattern.

---

As for your question, in comments, consult this snippet in the OpenCL specification, section 6.5:

// declares a pointer p in the __private address space that
// points to an int object in address space __global
__global int *p;

So the memory space associated with pointer types represents the memory space of the variable they point to, not the pointers themselves, which are always __private (i.e. per work item).

---

You cannot use string manipulation functions from the standard C library in OpenCL, though you can recode them for use on the GPU if you want to (most of them are hardly difficult to implement).