如何在OpenCL中对全局计数器进行原子递增

Question

我希望在OpenCL中拥有一个全局计数器，可以由每个工作组中的每个工作项增加。

在我的内核中，我这样做：

#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);
}

我的主机代码是使用最少的pyopencl来排队内核：

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()

我希望输出显示一个（可能是随机的）外观为"Counter：i"，其中i是总工作项的数量（在我的情况下为2）。但实际上我没有得到任何输出。当我重新运行程序时，它失败了。

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

彻底关闭我的IDE（Spyder）。

- Dschoni

Answer 1

volatile __global int* counter = 0; 

创建指向全局内存的指针是不够的，它后面必须有一些全局内存存储。

至少有两种选择：

1）如果你使用OpenCL 2.0实现，可以创建一个程序范围变量：

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) 创建一个 OpenCL 缓冲区并将其作为内核参数传递：

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);
}

- Bartosz Sochacki

---

看起来不错，但是1）对我不起作用，因为我不在OpenCL 2上；2）给我一个错误：“指针参数的指针参数的地址空间无效__kernel函数 __kernel void test(int counterArg)” 和 “从地址空间“private”到地址空间“global”的隐式转换在初始化表达式中不受支持 volatile __global int counterPtr = counterArg; ” - Dschoni

更改为__kernel void test(__global int *counterArg)似乎可以解决问题。 - Dschoni