global_work_size和local_work_size对应用程序逻辑有影响吗？

Question

我试图了解维度中所有不同的维度参数是如何组合在一起的.如果我的问题不明确,那部分是因为一个形成良好的问题需要我没有的答案.

work_dim,global_work_size和local_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的问题.你可以告诉我,我对所有这些都非常困惑,所以你能提供的任何帮助都会......帮助.

---

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

我在谷歌上发现的图像描述工作组

Answer 1

既然你曾表示自己对执行空间中涉及的概念有点困惑,我会在回答你的问题并举一些例子之前尝试总结它们.

线程/工作项在NDRange中组织,可以将其视为1,2,3个dims的网格.NDRange主要用于将每个线程映射到它们必须操纵的数据片段.因此,每个线程都应该唯一标识,并且线程应该知道它是哪一个以及它在NDRange中的位置.还有工作项内置函数.所有线程都可以调用这些函数来为它们提供有关它们的信息以及它们所处的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范围的第3列和第2行中的线程.该函数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)只有一个线程,但是有6个线程具有本地id(0,0),因为有6个组.

例:

这是一个将所有这些概念结合在一起的虚拟示例(请注意,性能会很糟糕,这只是一个愚蠢的例子).

假设您有一个包含6行和4列int的2D数组.您希望将这些元素分组为2乘2的元素,并以这样的方式对它们求和,例如,元素(0,0),(0,1),(1,0),(1,1)将在一个小组(希望它足够清楚).因为你将有6个"正方形",你将得到6个总和的结果,所以你需要一个包含6个元素的数组来存储这些结果.

要解决此问题,请使用上面详细介绍的2D NDRange.每个线程将从全局内存中获取一个元素,并将其存储在本地内存中.然后在同步之后,每个工作组只有一个线程,假设每个本地(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乘2平方).

另请注意,如果您决定使用尺寸为24且局部尺寸为4合1的1维NDRange,那么因为内核设计为使用2维,所以它也会搞砸.