全球和本地工作规模如何在此职能部门中分配？

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

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

如果 workgroup_size 确实是 4096，gsize 将最终成为 { 128, 32 }，因为它遵循 else 逻辑。 (>1024)
w 是 num_pixels_per_work_item = 32 宽列的数量，或者覆盖整个宽度的最小工作项数，对于 1920 的图像宽度是 60。换句话说，我们至少需要 60 x 1080 个工作项才能覆盖整个图像。
接下来计算分组列数和行数，暂存于global_work_size。由于组宽度已设置为 128，w 为 60 意味着我们最终得到 1 列组。（这似乎是一种资源浪费，每个组中的 128 个工作项中有一半以上不会做任何事情。）组行数只是 image_height 除以 gsize[1] (32) 和围捕。 (33.75 -> 34)
现在可以通过乘以网格来确定组的总数：num_groups = global_work_size[0] * global_work_size[1]
为了获得每个维度中工作项的真实总数，global_work_size 的每个维度现在乘以该维度中的组大小。 1, 34 乘以 128, 32 得到 128, 1088.

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

例如，如果我们反转 gsize[0] 和 gsize[1]，在这种情况下，我们将得到一个 { 32, 128 } 的组大小，从而得到一个 [=28] 的全局工作大小=] 并且只有 12% 的浪费。如果始终选择最大可能的组大小是一个好主意，那也值得检查；很可能不是，但我没有详细研究内核的计算，更不用说运行任何测量值，来判断是否是这种情况。

全球和本地工作规模如何在此职能部门中分配？

How are global and local work sizes distributed in this function?

parallel-processing

opencl