opencl - OpenCL 中减少总和的最佳工作组大小

Question

我正在使用以下内核进行求和缩减。

__kernel void reduce(__global float* input, __global float* output, __local float* sdata)
{
    // load shared mem
    unsigned int tid = get_local_id(0);
    unsigned int bid = get_group_id(0);
    unsigned int gid = get_global_id(0);

    unsigned int localSize = get_local_size(0);
    unsigned int stride = gid * 2;
    sdata[tid] = input[stride] + input[stride + 1];

    barrier(CLK_LOCAL_MEM_FENCE);
    // do reduction in shared mem
    for(unsigned int s = localSize >> 2; s > 0; s >>= 1) 
    {
        if(tid < s) 
        {
            sdata[tid] += sdata[tid + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // write result for this block to global mem
    if(tid == 0) output[bid] = sdata[0];
}

它工作正常，但如果我需要多个工作组(例如，如果我想计算 1048576 个元素的总和)，我不知道如何选择最佳工作组大小或工作组数量。据我了解，我使用的工作组越多，我得到的子结果就越多，这也意味着我最终需要更多的全局缩减。

我已经看到了一般工作组大小问题的答案 here 。有没有专门针对归约操作的建议？

Answer 1

这个问题可能与我不久前回答的问题重复: What is the algorithm to determine optimal work group size and number of workgroup .

实验将是确定任何给定设备的最佳方法。

更新: 我认为您可以安全地坚持一维工作组，就像您在示例代码中所做的那样。在主机上，您可以尝试最佳值。

对于每个设备:

1) 查询 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE。

2) 循环几个倍数并以该组大小运行内核。节省每个测试的执行时间。

3) 当您认为拥有最佳值时，请将其硬编码到新内核中以与该特定设备一起使用。这将进一步提升性能。您还可以消除设备特定内核中的 sdata 参数。

//define your own context, kernel, queue here

int err;
size_t global_size; //set this somewhere to match your test data size
size_t preferred_size;
size_t max_group_size;

err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), preferred_size, NULL);
//check err
err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), max_group_size, NULL);
//check err

size_t test_size;

//your vars for hi-res timer go here

for (unsigned int i=preferred_size ; i<=max_group_size ; i+=preferred_size){
    //reset timer
    test_size = (size_t)i;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &test_size, 0, NULL, NULL);
    if(err){
        fail("Unable to enqueue kernel");  //implement your own fail function somewhere..
    }else{
        clfinish(queue);
        //stop timer, save value
        //output timer value and test_size
    }
}

特定于设备的内核可能如下所示，但第一行应替换为您的最佳值:

#define LOCAL_SIZE 32
__kernel void reduce(__global float* input, __global float* output)
{
    unsigned int tid = get_local_id(0);
    unsigned int stride = get_global_id(0) * 2;
    __local float sdata[LOCAL_SIZE];
    sdata[tid] = input[stride] + input[stride + 1];

    barrier(CLK_LOCAL_MEM_FENCE);

    for(unsigned int s = LOCAL_SIZE >> 2; s > 0; s >>= 1){
        if(tid < s){
            sdata[tid] += sdata[tid + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if(tid == 0) output[get_group_id(0)] = sdata[0];
}