I would like to use something like a “static int” inside a kernel - for example like a counter that always increments each time the kernel is called.
But code like this does not compile:
__kernel void vectorAdd(__global const float * a,
__global float * c)
{
static int counter = 0;
int nIndex = get_global_id(0);
c[nIndex] = a[nIndex] + c[nIndex];
counter++;
}
Is there something equivalent to static in OpenCL?
The closest thing you can do is probably a __global kernel argument and use atomic operations to increment/decrement.
You surely realize that a kernel function is executed once for each work-item. If you execute and NDRange with 1000x1000 work-items the kernel is evaluated 1000000 times. Atomic operations are necessary to ensure that you get the correct results since multiple work-items may be running simultaneously.
Yes I do. Was just an example. In the end I’d like to write a random number generator and therefore I need a global variable.
Maybe that should have been my question: Is there something like a global variable in OpenCL. Preferably stored on the device - so there is no memory transfer host <-> device involved.
As for atomic. Would it be possible that you give me a short example?
cheers!
In OpenCL all variables at global scope are in the __constant address space. You need to add a new __global argument to all of your kernels if you want to have global mutable state.
As for atomics, search the spec for a section titled “Atomic Functions”. Let me know if it’s still confusing.
Thank you for your answer.
How/where do I define such a variable? From what I’ve seen so far everything in OpenCL is inside kernels/functions. Is it possible, like in C/C++, to define/declare something outside those kernels?
How/where do I define such a variable?
I think you may be a bit confused by the term “__global”. It doesn’t refer to the scope of the variable, but to the memory space where it is stored. Let me explain this with a couple of examples.
Variable “foo” is at global scope:
__constant float foo = 1.234;
__kernel void mykernel(...)
{
}
Variables at global scope (also called “program scope”) have to be declared in the __constant address space, which means that their value is immutable.
A completely different notion is “__global” variables. The __global here refers to the type of memory where they are stored. __global memory is shared by all work-items in all work-groups. Since everybody can access it, the user must be careful when it writes into it, because other work-items could be reading from or writing to the same memory location at the same time.
Here, “bar” is a kernel argument that points to __global memory:
__kernel void mykernel(__global float *bar)
{
}
You can use a __global kernel argument like this to store state that is shared across all work-items and across successive kernel invocations, which seems to be what you are trying to do.