在这个函数中，全局和局部工作规模如何分配？

Question

这是来自 样板戏 我对全局和局部工作尺寸是如何计算的感到困惑，它们是根据图像尺寸计算的。

图像尺寸是1920 x 1080 (w x h)。

我以为 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×1088。

设备上最大可用的group_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 落得 { 128, 32 } 由于它 else 逻。(>1024)
2. w 是多少 num_pixels_per_work_item = 32 宽栏，或覆盖整个宽度的最小工作项数，对于宽度为1920的图像来说是60。换句话说，我们要求覆盖整个图像的绝对最小工作项数为60×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%的浪费。这也是值得检查的，如果总是选择尽可能大的组大小是一个好主意；很可能不是，但我还没有详细研究内核的计算，更不用说运行任何测量，说是否是这样。