在 OpenCL 中减少的最佳实践是什么？

Question

想象一个具有关联性的二元运算（我们将其命名为“+”）属性。当你可以并行计算 a1 + a2 + a3 + a4 + ... 时，首先计算

b1 = a1 + a2
b2 = a3 + a4

然后

c1 = b1 + b2
c2 = b3 + b4

然后对上一步的结果做同样的事情，依此类推，直到剩下一个元素。

我正在学习 OpenCL 并尝试实施这种方法来汇总数组中的所有元素。我是这项技术的新手，所以该程序可能看起来有些奇怪。

这是内核：

__kernel void reduce (__global float *input, __global float *output)
{
    size_t gl = get_global_id (0);
    size_t s = get_local_size (0);
    int i;
    float accum = 0;

    for (i=0; i<s; i++) {
        accum += input[s*gl+i];
    }

    output[gl] = accum;
}

这是主程序：

#include <stdio.h>
#include <stdlib.h>
#include <fcntl.h>
#include <unistd.h>
#include <sys/mman.h>
#include <sys/stat.h>
#include <CL/cl.h>

#define N (64*64*64*64)

#include <sys/time.h>
#include <stdlib.h>

double gettime ()
{
    struct timeval tv;
    gettimeofday (&tv, NULL);
    return (double)tv.tv_sec + (0.000001 * (double)tv.tv_usec);
}

int main()
{
    int i, fd, res = 0;
    void* kernel_source = MAP_FAILED;

    cl_context context;
    cl_context_properties properties[3];
    cl_kernel kernel;
    cl_command_queue command_queue;
    cl_program program;
    cl_int err;
    cl_uint num_of_platforms=0;
    cl_platform_id platform_id;
    cl_device_id device_id;
    cl_uint num_of_devices=0;
    cl_mem input, output;
    size_t global, local;

    cl_float *array = malloc (sizeof (cl_float)*N);
    cl_float *array2 = malloc (sizeof (cl_float)*N);
    for (i=0; i<N; i++) array[i] = i;

    fd = open ("kernel.cl", O_RDONLY);
    if (fd == -1) {
        perror ("Cannot open kernel");
        res = 1;
        goto cleanup;
    }
    struct stat s;

    res = fstat (fd, &s);
    if (res == -1) {
        perror ("Cannot stat() kernel");
        res = 1;
        goto cleanup;
    }

    kernel_source = mmap (NULL, s.st_size, PROT_READ, MAP_PRIVATE, fd, 0);
    if (kernel_source == MAP_FAILED) {
        perror ("Cannot map() kernel");
        res = 1;
        goto cleanup;
    }

    if (clGetPlatformIDs (1, &platform_id, &num_of_platforms) != CL_SUCCESS) {
        printf("Unable to get platform_id\n");
        res = 1;
        goto cleanup;
    }

    if (clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id,
                       &num_of_devices) != CL_SUCCESS)
    { 
        printf("Unable to get device_id\n");
        res = 1;
        goto cleanup;
    }
    properties[0]= CL_CONTEXT_PLATFORM;
    properties[1]= (cl_context_properties) platform_id;
    properties[2]= 0;
    context = clCreateContext(properties,1,&device_id,NULL,NULL,&err);
    command_queue = clCreateCommandQueue(context, device_id, 0, &err);
    program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, NULL, &err);


    if (clBuildProgram(program, 0, NULL, NULL, NULL, NULL) != CL_SUCCESS) {
        char buffer[4096];
        size_t len;

        printf("Error building program\n");
        clGetProgramBuildInfo (program, device_id, CL_PROGRAM_BUILD_LOG, sizeof (buffer), buffer, &len);
        printf ("%s\n", buffer);
        res = 1;
        goto cleanup;
     }

    kernel = clCreateKernel(program, "reduce", &err);
    if (err != CL_SUCCESS) {
        printf("Unable to create kernel\n");
        res = 1;
        goto cleanup;
    }

    // create buffers for the input and ouput
    input = clCreateBuffer(context, CL_MEM_READ_ONLY, 
                            sizeof(cl_float) * N, NULL, NULL);
    output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
                            sizeof(cl_float) * N, NULL, NULL);

    // load data into the input buffer
    clEnqueueWriteBuffer(command_queue, input, CL_TRUE, 0, 
                          sizeof(cl_float) * N, array, 0, NULL, NULL);

    size_t size = N;
    cl_mem tmp;
    double time = gettime();
    while (size > 1)
    {
        // set the argument list for the kernel command
        clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
        clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
        global = size;
        local = 64;

        // enqueue the kernel command for execution
        clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
                           &local, 0, NULL, NULL);
        clFinish(command_queue);
        size = size/64;
        tmp = output;
        output = input;
        input = tmp;
    }
    cl_float answer[1];
    clEnqueueReadBuffer(command_queue, tmp, CL_TRUE, 0, 
                        sizeof(cl_float), array, 0, NULL, NULL);
    time = gettime() - time;
    printf ("%f %f\n", array[0], time);

cleanup:
    free (array);
    free (array2);
    clReleaseMemObject(input);
    clReleaseMemObject(output);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(command_queue);
    clReleaseContext(context);

    if (kernel_source != MAP_FAILED) munmap (kernel_source, s.st_size);
    if (fd != -1) close (fd);

    _Exit (res); // Kludge
    return res;
}

所以我重新运行内核，直到缓冲区中只有一个元素。这是计算 OpenCL 中元素总和的正确方法吗？我用 gettime 测量的时间比 CPU 上的简单循环的执行时间慢大约 10 倍（编译的 clang 4.0.0 和 -O2 -ffast-math 标志）。我使用的硬件：AMD Ryzen 5 1600X 和 Amd Radeon HD 6950。

Answer 1

您可以采取一些措施来提高性能。

首先，去掉循环中的 clFinish 调用。这迫使内核的单独执行依赖于命令队列的整个状态在继续之前达到与主机的同步点，这是不必要的。唯一需要的同步是内核按顺序执行，即使你有一个乱序队列（你的程序无论如何都不会请求），你可以通过简单地使用事件对象来保证这一点。

size_t size = N;
size_t total_expected_events = 0;
for(size_t event_count = size; event_count > 1; event_count /= 64)
    total_expected_events++;
cl_event * events = malloc(total_expected_events * sizeof(cl_event));
cl_mem tmp;
double time = gettime();
size_t event_index = 0;
while (size > 1)
{
    // set the argument list for the kernel command
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
    global = size;
    local = 64;

    if(event_index == 0)
        // enqueue the kernel command for execution
        clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
                           &local, 0, NULL, events);
    else
        clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
                           &local, 1, events + (event_index - 1), events + event_index);
    size = size/64;
    tmp = output;
    output = input;
    input = tmp;
    event_index++;
}
clFinish(command_queue);
for(; event_index > 0; event_index--)
    clReleaseEvent(events[event_index-1]);
free(events);
cl_float answer[1];
clEnqueueReadBuffer(command_queue, tmp, CL_TRUE, 0, 
                    sizeof(cl_float), array, 0, NULL, NULL);

另一件可能需要研究的事情是在一个内核中执行全部缩减，而不是将其分散到同一内核的多个调用中。 This is one potential 示例，尽管它可能比您需要的更复杂。

在 OpenCL 中减少的最佳实践是什么？

What is the best practice to do reduce in OpenCL?

reduce

opencl