c - OpenCL 内核似乎没有获取全局 id "globally"

Question

我正在尝试将我编写的程序转换为 OpenCL，但我对它还不够熟悉。尽管如此，我的(三个)内核之一还是遇到了麻烦。它基本上是一个复杂的矩阵 vector 乘法，但我编写它是为了更好地满足我的需求。

问题是，我无法让内核在 GPU 上工作。我已经把它简化到最多(2行)，在CPU上调试，并且它在CPU上运行得很好。但当涉及到 GPU 时，一切都搞砸了。我在 MacBook Pro 上工作，在 NVIDIA GeForce 650M 上得到一个结果，而在集成 Intel HD 4000 上得到另一个结果。内核是

__kernel void Chmv_(__global float2 *H, const float alpha, __global float2 *vec, 
                const int off/*in number of elements*/,
                __local float2 *vw,
                __global float2 *vout) 
{
int gidx=get_global_id(0);
int gidy=get_global_id(1);
int gs=get_global_size(0);

    vout[gidx].x += alpha*(H[gidx+gidy*gs].x*vec[gidy].x-H[gidx+gidy*gs].y*vec[gidy].y);
    vout[gidx].y += alpha*(H[gidx+gidy*gs].y*vec[gidy].x+H[gidx+gidy*gs].x*vec[gidy].y);

}

对于测试，我让矩阵 H 为 4x4 矩阵，填充为 (1.0f, 0.0f)，而输入 vector vec 具有 x 个分量 (0.0 、1.0、2.0、3.0) 和 y 分量 0。alpha 设置为 2.0f。所以，我应该有 (12, 12, 12, 12) 作为 x 输出，如果我使用 CPU，我就会这么做。 NVIDIA 给我 6.0，而 Intel 给我 4.0。

现在，仔细检查后发现，如果输入 vector 是 (0,1,2,0)，NVIDIA 会给出 0 作为答案，如果是 (0,1,0,3)，Intel 会给出 0 作为答案出色地。顺便说一句，将 vec[gidy] 更改为 vec[gidx] 使 vector 加倍。从这些来看，在我看来，内核仅在一维 x 上执行良好，而 get_global_id(1) 只有一个值，这显然是不行的。

我将添加调用此内核检查的测试函数。现在，有人知道会发生什么吗？

void _test_(){
cl_mem mat,vec, out;
size_t gs[2]={4,4};
size_t ls[2]={1,4};
size_t cpuws[2]={1,1};
cl_float2 *A=(cl_float2*)calloc(gs[0]*gs[0], sizeof(cl_float2));
cl_float2 *v=(cl_float2*)calloc(gs[0], sizeof(cl_float2));
cl_float2 *w=(cl_float2*)calloc(gs[0], sizeof(cl_float2));
int i;

for (i=0; i<gs[0]; i++) {
    A[i*gs[0]].x=1.0;
    A[i*gs[0]+1].x= 1.0;//(i<ls-1)? 1.0f:0.0f;
    A[i*gs[0]+2].x=1.0;
    A[i*gs[0]+3].x=1.0;
    v[i].x=  (float)i;
    printf("%d %f %f %f %f\n%v2f\n",i, A[i*gs[0]].x, A[i*gs[0]+1].x, A[i*gs[0]+2].x, A[i*gs[0]+3].x, v[i]);
}
v[2].x=0.0f; //<--- set individually for debug

mat = clCreateBuffer(context, CL_MEM_READ_WRITE, gs[0]*gs[0]*sizeof(cl_float2), NULL, NULL);
vec = clCreateBuffer(context, CL_MEM_READ_WRITE, gs[0]*sizeof(cl_float2), NULL, NULL);
out = clCreateBuffer(context, CL_MEM_READ_WRITE, gs[0]*sizeof(cl_float2), NULL, NULL);

error = clEnqueueWriteBuffer(queue, mat, CL_TRUE, 0, gs[0]*gs[0]*sizeof(cl_float2), A, 0, NULL, NULL);
error = clEnqueueWriteBuffer(queue, vec, CL_TRUE, 0, gs[0]*sizeof(cl_float2), v, 0, NULL, NULL);
error = clEnqueueWriteBuffer(queue, out, CL_TRUE, 0, gs[0]*sizeof(cl_float2), w, 0, NULL, NULL);

int offset=0;
float alpha=2.0;
error  = clSetKernelArg(Chmv_, 0, sizeof(cl_mem),&mat);
error |= clSetKernelArg(Chmv_, 1, sizeof(float), &alpha);
error |= clSetKernelArg(Chmv_, 2, sizeof(cl_mem),&vec);
error |= clSetKernelArg(Chmv_, 3, sizeof(int), &offset);
error |= clSetKernelArg(Chmv_, 4, gs[0]*sizeof(cl_float2), NULL);
error |= clSetKernelArg(Chmv_, 5, sizeof(cl_mem), &out);
assert(error == CL_SUCCESS);

error = clEnqueueNDRangeKernel(queue, Chmv_, 2, NULL, gs, NULL, 0, NULL, &event);

error = clEnqueueReadBuffer(queue, out, CL_TRUE, 0, gs[0]*sizeof(cl_float2), w, 0, NULL, NULL);
clFinish(queue);

for (i=0; i<gs[0]; i++) {
    printf("%f %f\n", w[i].x, w[i].y);

}

clReleaseMemObject(mat);
clReleaseMemObject(vec);
clReleaseMemObject(out);
}

Answer 1

您遇到了对公共(public)内存区域的多线程不安全访问的典型问题。 (vout)

您必须认为所有工作项都会同时运行。这意味着，它们将以任何顺序读取和写入内存。

当您在CPU中执行时，问题不会出现，因为执行是由硬件串行完成的。 然而在 GPU 中，一些工作项会读取 vout 的内存，对其进行递增并写入。但其他人也会在之前的工作项写入新值之前读取 vout 的内存。

可能您的所有工作项都在并行运行，因为您的内核大小很小，这就是为什么您只能看到其中一个添加到最终结果中。

这是一个典型的并行归约问题。你可以谷歌一下了解更多细节。您需要实现的是在访问 vout 时同步所有线程，可以通过 atomic_add() (慢)或通过适当的减少(很难编码)。你可以查看这个指南，它是针对 CUDA 的，但基本思想或多或少相同:Reduction Guide