OpenCL 内核似乎没有“全局”获取全局 ID

Question

我正在尝试将我所做的程序转换为 OpenCL，但我对它还不够熟悉。尽管如此，我的（三个）内核之一还是有问题。它基本上是一个复杂的矩阵向量乘法，但我写它是为了更好地满足我的需要。

问题是，我无法让内核在 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);

}

对于测试，我让 Matrix H 是一个 4x4 矩阵，用 (1.0f, 0.0f) 填充，而输入向量 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。

现在，仔细检查表明，如果输入向量是 (0,1,2,0)，NVIDIA 给我 0 作为答案，如果是 (0,1,0,3)，英特尔给我 0 作为答案好吧。顺便说一句，将vec[gidy] 更改为vec[gidx] 只会使向量翻倍。从这些来看，在我看来，内核只在一维 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

您遇到了一个典型的问题，即多线程不安全地访问公共内存区域。 (vout)

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

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

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

这是一个典型的并行约简问题。您可以谷歌它以获取更多详细信息。您需要实现的是在访问vout 时同步所有线程，无论是通过atomic_add()（慢）还是通过适当的减少（难以编码）。您可以查看本指南，它适用于 CUDA，但基本思想大致相同：Reduction Guide