由于您表示自己对执行空间中涉及的概念有些困惑,因此在回答您的问题并给出一些示例之前,我将尝试对其进行总结。
线程/工作项组织在 NDRange 中,可以将其视为 1、2、3 暗度的网格。NDRange 主要用于将每个线程映射到每个线程都必须操作的数据。因此,每个线程都应该被唯一标识,并且一个线程应该知道它是哪一个以及它在 NDRange 中的位置。还有 Work-Item 内置函数。所有线程都可以调用这些函数,以向它们提供有关它们自己和它们所在的 NDRange 的信息。
尺寸:
如前所述,NDRange 最多可以有 3 个维度。因此,如果您以这种方式设置尺寸:
size_t global_work_size[2] = { 4, 4 };
这并不意味着每个维度将有 4 个工作组,总共 8 个,而是您的 NDRange 中将有 4 * 4 即 16 个线程。这些线将排列成一个“正方形”,边长为 4 个单位。uint get_work_dim ()
使用该函数,工作项可以知道 NDRange 由多少个维度组成。
全球规模:
线程还可以使用 查询特定维度的 NDRange 有多大size_t get_global_size (uint D)
。因此他们可以知道“线/正方形/矩形/立方体”NDRange有多大。
全局唯一标识符:
由于该组织,每个线程都可以使用与特定维度相对应的索引来唯一标识。因此,线程 (2, 1) 指的是位于 2D 范围的第三列和第二行的线程。该函数size_t get_global_id (uint D)
在内核中用于查询线程的id。
工作组(或本地)大小:
NDRange 可以分成更小的组,称为工作组。这是您所指的 local_work_size 也(并且在逻辑上)最多 3 个维度。请注意,对于低于 2.0 的 OpenCL 版本,给定维度中的 NDRange 大小必须是该维度中工作组大小的倍数。所以为了保持你的例子,因为在维度 0 中我们有 4 个线程,所以维度 0 中的工作组大小可以是 1、2、4 但不是 3。与全局大小类似,线程可以使用 查询本地大小size_t get_local_size (uint D)
。
本地唯一标识符:
有时,可以在工作组中唯一标识线程很重要。因此函数size_t get_local_id (uint D)
。注意上句中的“内”。具有本地 id (1, 0) 的线程将是唯一在其工作组(2D)中具有此 id 的线程。但是,具有本地 id (1, 0) 的线程将与 NDRange 中的工作组一样多。
组数:
有时说到组,一个线程可能需要知道有多少组。这就是函数size_t get_num_groups (uint D)
存在的原因。请注意,您必须再次将您感兴趣的维度作为参数传递。
每个组还有一个 id:
...您可以使用函数在内核中查询size_t get_group_id (uint D)
。请注意,组 ID 的格式将类似于线程的格式:最多 3 个元素的元组。
概括:
总结一下,如果您有一个全局工作大小为 (4, 6) 和局部工作大小为 (2, 2) 的2D NDRange,这意味着:
- 维度 0 中的全局大小将为 4
- 维度 1 中的全局大小将为 6
- 维度 0 中的本地大小(或工作组大小)将为 2
- 维度 1 中的本地大小(或工作组大小)将为 2
- 维度 0 中的线程全局 ID 范围为 0 到 3
- 维度 1 中的线程全局 ID 范围为 0 到 5
- 维度 0 中的线程本地 ID 范围为 0 到 1
- 维度 1 中的线程本地 ID 范围为 0 到 1
- NDRange 中的线程总数将为 4 * 6 = 24
- 工作组中的线程总数将为 2 * 2 = 4
- 工作组总数将为 (4/2) * (6/2) = 6
- 维度 0 中的组 ID 范围为 0 到 1
- 维度 1 中的组 ID 范围为 0 到 2
- 全局 id (0, 0) 只有一个线程,但本地 id (0, 0) 将有 6 个线程,因为有 6 个组。
例子:
这是一个将所有这些概念一起使用的虚拟示例(请注意,性能会很糟糕,这只是一个愚蠢的示例)。
假设您有一个 6 行和 4 列 int 的二维数组。您希望将这些元素按 2 x 2 个元素的平方分组,并以这样的方式对它们进行汇总,例如,元素 (0, 0)、(0, 1)、(1, 0)、(1, 1)将在一个组中(希望它足够清楚)。因为您将有 6 个“正方形”,所以您将有 6 个求和结果,因此您需要一个包含 6 个元素的数组来存储这些结果。
为了解决这个问题,您可以使用上面详述的 2D NDRange。每个线程将从全局内存中获取一个元素,并将其存储在本地内存中。然后在同步之后,每个工作组只有一个线程,假设每个 local(0, 0) 线程将对元素(本地)求和,然后将结果存储在 6 元素数组(全局)中的特定位置。
//in is a 24 int array, result is a 6 int array, temp is a 4 int array
kernel void foo(global int *in, global int *result, local int *temp){
//use vectors for conciseness
int2 globalId = (int2)(get_global_id(0), get_global_id(1));
int2 localId = (int2)(get_local_id(0), get_local_id(1));
int2 groupId = (int2)(get_group_id (0), get_group_id (1));
int2 globalSize = (int2)(get_global_size(0), get_global_size(1));
int2 locallSize = (int2)(get_local_size(0), get_local_size(1));
int2 numberOfGrp = (int2)(get_num_groups (0), get_num_groups (1));
//Read from global and store to local
temp[localId.x + localId.y * localSize.x] = in[globalId.x + globalId.y * globalSize.x];
//Sync
barrier(CLK_LOCAL_MEM_FENCE);
//Only the threads with local id (0, 0) sum elements up
if(localId.x == 0 && localId.y == 0){
int sum = 0;
for(int i = 0; i < locallSize.x * locallSize.y ; i++){
sum += temp[i];
}
//store result in global
result[groupId.x + numberOfGrp.x * groupId.y] = sum;
}
}
最后回答您的问题: global_work_size 和 local_work_size 对应用程序逻辑有影响吗?
通常是的,因为它是您设计算法的方式的一部分。请注意,工作组的大小不是随机取的,而是符合我的需要(这里是 2 x 2 平方)。
另请注意,如果您决定使用尺寸为 24 的 NDRange 尺寸为 24 且局部尺寸为 4 in 1 dim,它也会搞砸,因为内核被设计为使用 2 尺寸。