代码之家 › 专栏 › 技术社区 › VHristov

OpenCL问题

opencl memory-management

VHristov · 技术社区 · 16 年前

__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
...

大约3200个元素之后,我得到了这个由80个元素组成的块。位置不总是一样的,但元素的数量总是一样的-80。更奇怪的是-如果我把第一行改成80 问题消失了。但是,在使用线程id之后,我得出的结论是,问题并没有消失,它只是移动了。使用线程1425,我有一半的时间得到问题,当我得到它时,buggy块在数组的末尾。因此,我假设,当我没有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个元素的块被另一个“线程”覆盖。但我想得越多,就越没有意义。

编辑:

#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的问题,它们又是0和1。此外,对于大小为1的全局工作组,它们都可以正常工作。也就是说,如果我使用两个内核,一个是全局工作组大小为[problem size]的并行计算,另一个是全局工作组大小为1的并行计算,一切都很好。再说一次,我很清楚这不是实现它的方法(一个运行这种顺序代码的内核),但是,我想知道为什么它不工作,因为看起来我遗漏了一些东西。

谢谢,

1 回复 | 直到 16 年前

Josep 16 年前

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

编辑: 在看过你的代码并进行测试之后,我意识到一开始我并没有完全理解你的问题。当你出现在面具更新部分时,我的脑子里就摆脱了那条线。我第一次应该能答对的。

问题是不能同步不同的工作组。CLK\u GLOBAL\u MEM\u FENCE影响工作组的内存顺序访问(确保在回读之前完成对全局内存的写入)。真正的解决方案是在两个调用中执行代码,首先并行更新掩码,然后在另一个内核中执行其余的工作,第一个调用完成后将执行这些工作。您需要在继续之前完成整个操作,因此必须在命令队列级别使用屏障。没有别的办法了。

规范中的逐字:

单个工作组中的工作项

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