全局和本地工作量如何在此功能中分布？

Question

来自a sample program，用于 OpenCL 编程。 我对如何计算全局和本地工作量感到困惑。 它们是根据图像大小计算的。

图像尺寸为 1920 x 1080（宽 x 高）。

我假设 global_work_size[0] 和 global_work_size[1] 是图像上的网格。

但现在 global_work_size 是 {128, 1088}。

那么 local_work_size[0] 和 local_work_size[1] 是 global_work_size 上的网格。 local_work_size 是 {128, 32}。

但总组数，num_groups = 34，不是 128 x 1088。

设备上可用的最大 workgroup_size 为 4096。

图像是如何分配到这样的全局和本地工作组大小的？

它们在以下函数中计算。

    clGetKernelWorkGroupInfo(histogram_rgba_unorm8, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
    {
        size_t  gsize[2];
        int     w;

        if (workgroup_size <= 256)
        {
            gsize[0] = 16;//workgroup_size is formed into row & col
            gsize[1] = workgroup_size / 16;
        }
        else if (workgroup_size <= 1024)
        {
            gsize[0] = workgroup_size / 16;
            gsize[1] = 16;
        }
        else
        {
            gsize[0] = workgroup_size / 32;
            gsize[1] = 32;
        }

        local_work_size[0] = gsize[0];
        local_work_size[1] = gsize[1];

        w = (image_width + num_pixels_per_work_item - 1) / num_pixels_per_work_item;//to include all pixels, num_pixels_per_work_item is added first
        global_work_size[0] = ((w + gsize[0] - 1) / gsize[0]);//col
        global_work_size[1] = ((image_height + gsize[1] - 1) / gsize[1]);//row

        num_groups = global_work_size[0] * global_work_size[1];    
        global_work_size[0] *= gsize[0];
        global_work_size[1] *= gsize[1];
    }    
    err = clEnqueueNDRangeKernel(queue, histogram_rgba_unorm8, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
    if (err)
    {
        printf("clEnqueueNDRangeKernel() failed for histogram_rgba_unorm8 kernel. (%d)\n", err);
        return EXIT_FAILURE;
    } 

Answer 1

我看不出这里有什么大秘密。如果您按照计算，这些值确实会如您所说的那样结束。 （在我看来，小组规模并不是特别有效。）

1. 如果 workgroup_size 确实是 4096，则 gsize 将按照 else 的逻辑结束 { 128, 32 }。 (>1024)
2. w 是 num_pixels_per_work_item = 32 宽列的数量，或覆盖整个宽度的最小工作项数，对于 1920 的图像宽度，它是 60。换句话说，我们要求绝对最小值为 60 x 1080覆盖整个图像的工作项。
3. 接下来，计算组列数和行数并将其临时存储在global_work_size 中。由于组宽度已设置为 128，w 为 60 意味着我们最终得到 1 列组。 （这似乎是一种资源浪费，每个组中的 128 个工作项中有一半以上不会做任何事情。）组行数只是 image_height 除以 gsize[1] (32) 并四舍五入。 (33.75 -> 34)
4. 现在可以通过乘以网格来确定组的总数：num_groups = global_work_size[0] * global_work_size[1]
5. 为了获得每个维度中工作项的真实总数，global_work_size 的每个维度现在乘以该维度中的组大小。 1, 34 乘以 128, 32 得到 128, 1088。

这实际上覆盖了 4096 x 1088 像素的区域，因此其中大约 53% 是浪费的。这主要是因为组维度的算法有利于宽组，并且每个工作项都在图像的 32x1 像素切片上工作。最好支持高大的工作组以减少四舍五入的数量。

例如，如果我们颠倒 gsize[0] 和 gsize[1]，在这种情况下，我们将得到 { 32, 128 } 的组大小，从而使我们的全局工作大小为 { 64, 1152 }，并且只有 12% 的浪费。如果总是选择尽可能大的群体规模是否是一个好主意，这也是值得检查的；很可能不是，但我没有详细研究过内核的计算，更不用说运行任何测量，以确定是否是这种情况。