1

我遇到了 OpenCL 的问题，我希望有人能提示可能是什么原因。以下是该程序的一个版本，减少了问题。我有一个大小为 4000 的输入 int 数组。在我的内核中，我正在进行扫描。显然，有很好的方法可以并行执行此操作，但要重现该问题，只有一个线程在执行整个计算。在扫描之前，输入 (result_mask) 只有值 0 或 1。

__kernel void
sel_a(__global db_tuple * input,
      __global int * result_mask,
      __global int * result_count,
      const unsigned int max_id)
{
// update mask based on input in parallel

mem_fence(CLK_GLOBAL_MEM_FENCE);

if(gid == 0)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}
}

预期的结果将是最初具有不同于 0 的值且结果掩码中只有 5 的元素的数量。然而，事实并非如此。输出如下所示：

...
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 
5 5 5 5 5 5 5 5 5 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 1 0 0 0 1 0 0 0 1 0 0 0 0 0 0 0 0 0 0 
0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 5 5 5 5 5 5 5 5 5 5 5 
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5
...

大约在某个地方后，我得到了这个 80 个元素的块。3200 个元素。它并不总是相同的位置，但它总是相同数量的元素 - 80。它变得更加奇怪 - 如果我将第一行更改为 if(gid == 2000) 问题就消失了。但是，在玩弄了线程 id 之后，我得出的结论是问题并没有消失，它只是移动了。使用线程 1425，我得到了一半的问题，当我得到它时，错误块位于数组的末尾。因此我假设，当我没有 0 和 1 时，该块已经“移动”得更远了。更令人兴奋的是——当我将输入大小增加到 5000 时，输出完全由 0 组成。此外，以下代码将不起作用：

if(gid == 0)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}
if(gid == 3999)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}

而只有

if(gid == 3999)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}

将起作用（同样，可能有更大的输入，它可能不起作用）。以下是该设备的一些详细信息：

Device name: GeForce 9600M GT
Device vendor: NVIDIA
    Clock frequency:        1250 MHz
    Max compute units:      4
    Global memory size:     256 MB
    Local memory size:.     16 KB
    Max memory allocation size: 128 MB
    Max work group size:        512

显然，我在这里错过了一些大事。我的第一个想法是这是一些内存冲突，其中 80 个元素的块被另一个“线程”覆盖。但我想得越多，它就越没有意义。

我将非常感谢任何提示！谢谢。

编辑：很抱歉回复晚了。所以我修改了代码，将其减少到最低限度以重现问题。以下是程序的c代码：

#include <stdio.h>
#include <stdlib.h>

#include <OpenCL/openCL.h>

#define INPUTSIZE (200)

typedef struct tag_openCL
{
    cl_device_id        device;

    cl_context          ctx;
    cl_command_queue    queue;
    cl_program          program;
} openCL;

int main(void)
{
    int err;
    openCL* cl_ctx = malloc(sizeof(openCL));

    if(!cl_ctx)
        exit(1);

    err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &cl_ctx->device, NULL);

    cl_ctx->ctx = clCreateContext(0, 1, &cl_ctx->device, clLogMessagesToStdoutAPPLE, NULL, &err);

    cl_ctx->queue = clCreateCommandQueue(cl_ctx->ctx, cl_ctx->device, CL_QUEUE_PROFILING_ENABLE, &err);

    printf("Successfully created context and queue for openCL device. \n");

    /* Build program */

    char * kernel_source = "__kernel void \
sel(__global int * input, \
    __global int * result_mask, \
    const unsigned int max_id) \
{ \
    int gid = get_global_id(0); \
    \
    result_mask[gid] = input[gid] % 2 == 0; \
    result_mask[gid] &= (input[gid] + 1) % 3 == 0; \
    \
    if(gid == 0) { \
        int i; \
        for(i = 0; i < max_id; i++) { \
            if(result_mask[i]) { \
                result_mask[i] = 5; \
            } \
            else { \
                result_mask[i] = 5; \
            } \
        } \
    } \
}";

    cl_program prog = clCreateProgramWithSource(cl_ctx->ctx, 1, (const char**)&kernel_source, NULL, &err);
    cl_ctx->program = prog;

    err = clBuildProgram(cl_ctx->program, 0, NULL, NULL, NULL, NULL);

    cl_kernel kernel = clCreateKernel(cl_ctx->program, "sel", &err);

    /* create dummy input data */
    int * input = calloc(sizeof(int), INPUTSIZE);
    int k;
    for(k = 0; k < INPUTSIZE; k++)
    {
        input[k] = abs((k % 5) - (k % 3))+ k % 2;
    }

    cl_mem source, intermediate;

    unsigned int problem_size = INPUTSIZE;

    source = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);
    clEnqueueWriteBuffer(cl_ctx->queue, source, CL_TRUE, 0, problem_size * sizeof(int), (void*) input, 0, NULL, NULL);

    intermediate = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);

    int arg = 0;
    clSetKernelArg(kernel, arg++, sizeof(cl_mem), &source);
    clSetKernelArg(kernel, arg++, sizeof(cl_mem), &intermediate);
    clSetKernelArg(kernel, arg++, sizeof(unsigned int), &problem_size);

    size_t global_work_size = problem_size;
    size_t local_work_size = 1;
    clEnqueueNDRangeKernel(cl_ctx->queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

    clFinish(cl_ctx->queue);

    // read results
    int * result = calloc(sizeof(int), problem_size );
    clEnqueueReadBuffer(cl_ctx->queue, intermediate, CL_TRUE, 0, problem_size * sizeof(int), result, 0, NULL, NULL);
    clFinish(cl_ctx->queue);


    int j;
    for(j=1; j<=problem_size; j++)
    {
        printf("%i \t", result[j-1]);
        if(j%10 ==0 && j>0)
            printf("\n");
    }

    return EXIT_SUCCESS;
}

结果仍然是不确定的，我在输出中的随机位置得到 0 和 1。对于大小为 1 的本地工作组，它们位于数组的前半部分，大小为 2 - 在后半部分，大小为 4 的 200 个元素看起来不错，但对于一个问题大小为 400。此外，对于 1 的全局工作组大小，一切正常。也就是说，如果我使用两个内核 - 一个执行全局工作组大小为 [问题大小] 的并行计算，另一个执行全局工作组大小为 1 的并行计算，一切都很好。同样，我完全清楚这不是这样做的方法（运行这种顺序代码的内核），但是，我想知道为什么它不起作用，因为它看起来我错过了一些东西。

谢谢，瓦西尔

4

1 回答 1

1

您的 OpenCL 代码非常简单，结果非常奇怪。我认为问题可能来自设置部分。缓冲区创建，调用 EnqueueNDRange 等。你能发布设置部分吗？我想问题可能在那里。

编辑：在看到您的代码并对其进行测试后，我意识到起初我并没有完全理解您的问题。当你评论面具更新部分时，我的头脑刚刚摆脱了那条线。我第一次应该能够正确回答。

问题是你不能同步不同的工作组。CLK_GLOBAL_MEM_FENCE 影响工作组的内存排序访问（确保在读回之前完成对全局内存的写入）。问题的真正解决方案是在两次调用中执行代码，首先并行更新掩码，然后在另一个内核中执行其余的工作，这些工作将在第一个调用完成时执行。您需要在继续之前完成整个操作，因此您必须在命令队列级别使用障碍。没有其他办法。

规范中的逐字记录：

OpenCL 中有两个同步域：

  • 单个工作组中的工作项

  • 在单个上下文中排队到命令队列的命令

于 2010-06-23T12:35:12.973 回答