这是来自 样板戏 我对全局和局部工作尺寸是如何计算的感到困惑,它们是根据图像尺寸计算的。
图像尺寸是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;
}
我看不出这里有什么大的奥秘。如果你按照计算,这些值确实如你所说。(在我看来,并不是说这组尺寸特别有效)。
workgroup_size
确实是4096。gsize
落得 { 128, 32 }
由于它 else
逻。(>1024)w
是多少 num_pixels_per_work_item = 32
宽栏,或覆盖整个宽度的最小工作项数,对于宽度为1920的图像来说是60。换句话说,我们要求覆盖整个图像的绝对最小工作项数为60×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 }
因此,我们的全球工作规模为 { 64, 1152 }
而只有12%的浪费。这也是值得检查的,如果总是选择尽可能大的组大小是一个好主意;很可能不是,但我还没有详细研究内核的计算,更不用说运行任何测量,说是否是这样。