10

我试图了解尺寸的所有不同参数如何在 OpenCL 中组合在一起。如果我的问题不清楚,部分原因是一个格式良好的问题需要一些我没有的答案。

work_dimglobal_work_sizelocal_work_size如何协同工作以创建您在内核中使用的执行空间?例如,如果我让 work_dim 2 那么我可以

get_global_id(0);
get_global_id(1);

我可以使用global_work_size将这两个维度划分为 n 个工作组,对吗?所以如果我像这样制作global_work_size

size_t global_work_size[] = { 4 };

那么每个维度会有 4 个工作组,总共 8 个?但是,作为一个初学者,我只使用 global_id 作为我的索引,所以无论如何只有全局 id 很重要。如您所知,我对所有这些都感到很困惑,因此您可以提供的任何帮助都将...帮助。


我制作的图像试图理解这个问题

我在谷歌上找到的描述工作组的图片

4

1 回答 1

41

由于您表示自己对执行空间中涉及的概念有些困惑,因此在回答您的问题并给出一些示例之前,我将尝试对其进行总结。

线程/工作项组织在 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 尺寸。

于 2013-08-14T10:48:09.847 回答