c++ - 在 OpenCL 中获得最佳的本地/全局工作组规模？

Question

我正在使用以下函数为我的 OpenCL 应用程序获取最佳的本地和工作组大小。

//maxWGSize == CL_KERNEL_WORK_GROUP_SIZE
//wgMultiple == CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
//compUnits == CL_DEVICE_MAX_COMPUTE_UNITS
//rems == max required work items

void MyOpenCL::getBestWGSize(cl_uint maxWGSize, cl_uint wgMultiple, cl_uint compUnits, cl_uint rems, size_t *gsize, size_t *lsize) const
{
    cl_uint cu = 1;
    if(wgMultiple <= rems)
    {
        bool flag = true;
        while(flag)
        {
            if(cu < compUnits)
            {
                cu++;
                if((wgMultiple * cu) > rems)
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
            else if(wgMultiple < maxWGSize)
            {
                wgMultiple *= 2;
                if((wgMultiple * cu) > rems)
                {
                    wgMultiple /= 2;
                    flag = false;
                    break;
                }
            }
            else
            {
                cu++;
                if(((wgMultiple * cu) > rems) || (cu > 2 * compUnits))
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
        }
    }
    else
    {
        bool flag = true;
        wgMultiple = 2;
        while(flag)
        {
            if(cu < compUnits)
            {
                cu++;
                if((wgMultiple * cu) > rems)
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
            else
            {
                wgMultiple *= 2;
                if((wgMultiple * cu) > rems)
                {
                    wgMultiple /= 2;
                    flag = false;
                    break;
                }
                else
                {
                    cl_int temp = rems - (wgMultiple * cu);
                    if((temp == 0) || (temp == 1))
                    {
                       flag = false;
                       break;
                    }
                }
            }
        }
    }

    *gsize = wgMultiple * cu;
    *lsize = wgMultiple;
    if(rems < *gsize)
    {
        *gsize = rems;
        *lsize = rems;
    }
    if(cu != compUnits)
    {
        while((cu * 2) <= compUnits)
        {
            cu *= 2;
            if(*lsize % 2 == 0)
                *lsize /= 2;
        }
    }
}

算法是:

1. 如果本地大小 ==，则决定需要多少个工作组CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
2. 如果还需要更多工作单元，则将局部大小乘以 2，直到它达到 CL_KERNEL_WORK_GROUP_SIZE

有什么改进算法的建议吗？

我得到的一些结果:

for GPU if max required work items == 99
maxWGSize    256 
wgMultiple   64 
compUnits    6 
rems     99 
*gsize   64 
*lsize   16 


for GPU if max required work items == 35
maxWGSize    256 
wgMultiple   4 
compUnits    6 
rems     35 
*gsize   24 
*lsize   4 

for GPU if max required work items == 57
maxWGSize    256 
wgMultiple   8 
compUnits    6 
rems     57 
*gsize   48 
*lsize   8 

for CPU if max required work items == 99
maxWGSize    1024 
wgMultiple   16 
compUnits    4 
rems     99 
*gsize   64 
*lsize   16 

for CPU if max required work items == 35
maxWGSize    1024 
wgMultiple   8 
compUnits    4 
rems     35 
*gsize   32 
*lsize   8

for CPU if max required work items == 57
maxWGSize    1024 
wgMultiple   8 
compUnits    4 
rems     57 
*gsize   32 
*lsize   8 

Answer 1

无可否认，我不明白(也几乎试图去理解)你在那里试图计算什么，因为它看起来过于复杂:确定最佳工作组规模应该几乎与计算单元的数量，应该没有必要用这么复杂的方式计算。

正如我在 original question 的回答中所说(并且由 DarkZeros in his comment 确认:只要您不使用本地内存等，您通常可以只传递 null 作为本地工作大小，OpenCL 会适本地选择它。

不过，可能会有一些注意事项。根据全局工作大小，底层 OpenCL 实现可能无法使用“良好”的本地工作组大小。例如:当全局工作大小为质数(大于最大本地工作大小)时，OpenCL 实现可能被迫使用本地工作大小 1...

这通常可以通过填充数据来规避，使其成为更合适的本地工作大小的倍数。首先，这意味着您必须修改内核，使其遵守工作大小的限制。在另一个问题的内核中，您必须为大小添加另一个参数，并相应地进行检查:

__kernel void reduceURatios(
    __global myreal *coef, 
    __global myreal *row, 
    myreal ratio,
    int sizeOfArrays)  // Add this parameter
{
    size_t gid = get_global_id(0);
    if (gid >= sizeOfArrays)
    {
        return; // Don't access invalid elements
    }

    myreal pCoef = coef[gid];
    myreal pRow = row[gid];

    pCoef = pCoef - (pRow * ratio);
    coef[gid] = pCoef;
}

然后您可以更自由地选择全局工作大小。当前问题的代码涉及 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE，但这与标准 GPU 上的这种微不足道的内核几乎无关。与此相反，CL_DEVICE_MAX_WORK_GROUP_SIZE 将是本地工作大小的不错选择(只要内核本身没有其他限制，例如寄存器压力 - 但这也绝对不是这里的情况)。

因此您可以使用 CL_DEVICE_MAX_WORK_GROUP_SIZE 作为计算全局工作大小的基础:

// As queried with CL_DEVICE_MAX_WORK_GROUP_SIZE
int maxWorkGroupSize = ...
int numWorkGroups = (n-1) / maxWorkGroupSize + 1;
int globalSizePadded = numWorkGroups * maxWorkGroupSize;

然后用这个(填充的)全局工作大小调用你的内核。您在内核中添加的 if 语句将确保线程不会访问无效的内存区域。当你用这个填充的全局大小启动你的内核，并将局部大小设置为 null 时，它应该自动选择 CL_DEVICE_MAX_WORK_GROUP_SIZE 作为局部大小(当然，你也可以手动指定)。

这可能使原始问题的计算速度更快，但它仍然不太可能比 CPU 版本更快...