static int inside a __kernel

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.