How to atomic increment a global counter in OpenCL

Question

I want to have a global counter in OpenCL that can be increased by every Work Item in every Work Group.

In my Kernel I do:

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

void increase(volatile __global int* counter)
{
    atomic_inc(counter);
}

__kernel void test()
{  
    volatile __global int* counter = 0; 
    increase(counter);
    printf("Counter: %i",&counter);
}

My Hostcode is minimum pyopencl to enqueue the Kernel:

import pyopencl as cl

platform = cl.get_platforms()[0]
devs = platform.get_devices()
device = devs[0]
ctx = cl.Context([device])
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags

f = open('Minimal.cl', 'r')
fstr = "".join(f.readlines())
prg = cl.Program(ctx, fstr).build()
test_knl = prg.test

def f():
     cl.enqueue_nd_range_kernel(queue,test_knl,(1,1,2),None)
f()

I expect an output to show a (possibly random ordered) Appearance of "Counter: i", where i is the number of total work items (in my case 2). Instead I don't get a printout. When I rerun the programm, it fails with

pyopencl.cffi_cl.LogicError: clcreatecontext failed: <unknown error -9999>

killing my IDE (Spyder) completely.

Answer 1

volatile __global int* counter = 0; 

Creating a pointer to global memory is not enough. There must to be some global memory storage behind it.

There are two options at least:

1) You can create a program scope variable if you use OpenCL 2.0 implementation:

void increase(volatile __global int* counter)
{
  atomic_inc(counter);
}

__global int counter = 0;

__kernel void test()
{  
  volatile __global int* counterPtr = &counter; 
  increase(counterPtr); // or increase(&counter);
  printf("Counter: %i",*counterPtr);
}

2) Create an OpenCL buffer and pass it by a kernel argument:

void increase(volatile __global int* counter)
{
  atomic_inc(counter);
}

__kernel void test(__global int *counterArg)
{  
  volatile __global int* counterPtr = counterArg; 
  increase(counterPtr); // or increase(counterArg);
  printf("Counter: %i",*counterPtr);
}