loops - OpenCL for 循环执行模型

Question

我目前正在学习 OpenCL 并遇到了这个代码片段:

int gti = get_global_id(0);
int ti = get_local_id(0);

int n = get_global_size(0);
int nt = get_local_size(0);
int nb = n/nt;

for(int jb=0; jb < nb; jb++) { /* Foreach block ... */
      pblock[ti] = pos_old[jb*nt+ti]; /* Cache ONE particle position */
      barrier(CLK_LOCAL_MEM_FENCE); /* Wait for others in the work-group */

      for(int j=0; j<nt; j++) { /* For ALL cached particle positions ... */
         float4 p2 = pblock[j]; /* Read a cached particle position */
         float4 d = p2 - p;
         float invr = rsqrt(d.x*d.x + d.y*d.y + d.z*d.z + eps);
         float f = p2.w*invr*invr*invr;
         a += f*d; /* Accumulate acceleration */
      }

      barrier(CLK_LOCAL_MEM_FENCE); /* Wait for others in work-group */
}

代码背景信息:这是 NBody 模拟程序中 OpenCL 内核的一部分。完整的代码和教程可以找到 here .

这是我的问题(主要与 for 循环有关):

OpenCL 中的 for 循环究竟是如何执行的？我知道所有工作项都运行相同的代码，并且工作组中的工作项尝试并行执行。因此，如果我在 OpenCL 中运行一个 for 循环，这是否意味着所有工作项都运行同一个循环，或者循环是否以某种方式划分为跨多个工作项运行，每个工作项执行循环的一部分(即工作项1 处理索引 0 ~ 9，项目 2 处理索引 10 ~ 19 等)。

在这段代码片段中，外循环和内循环是如何执行的？ OpenCL 是否知道外循环在所有工作组之间划分工作，而内循环试图在每个工作组内的工作项之间划分工作？

如果内部循环在工作项之间进行划分(意味着 for 循环中的代码是并行执行的，或者至少尝试并行执行)，那么最后的加法是如何工作的？它本质上是在做 a = a + f*d，根据我对流水线处理器的理解，这必须按顺序执行。

我希望我的问题足够清楚，我很感激任何意见。

Answer 1

1) How exactly are for-loops executed in OpenCL? I know that all work-items run the same code and that work-items within a work group tries to execute in parallel. So if I run a for loop in OpenCL, does that mean all work-items run the same loop or is the loop somehow divided up to run across multiple work items, with each work item executing a part of the loop (ie. work item 1 processes indices 0 ~ 9, item 2 processes indices 10 ~ 19, etc).

你说的对。所有工作项都运行相同的代码，但请注意，它们可能不会以相同的速度运行相同的代码。只有在逻辑上，它们运行相同的代码。在硬件中，同一wave(AMD术语)或warp(NV术语)内的工作项，它们完全遵循指令级别的足迹。

就循环而言，无非就是汇编代码级别的几个分支操作。来自同一波的线程并行执行分支指令。如果所有工作项都满足相同的条件，那么它们仍然遵循相同的路径，并并行运行。但是，如果他们不同意相同的条件，那么通常会出现不同的执行。例如，在下面的代码中:

if(condition is true)
   do_a();
else
   do_b();

从逻辑上讲，如果某些工作项满足条件，它们将执行 do_a() 函数；而其他工作项将执行 do_b() 函数。然而，实际上，一个wave中的工作项在硬件中执行完全相同的步骤，因此它们不可能并行运行不同的代码。因此，一些工作项将被 do_a() 操作屏蔽掉，而 wave 执行 do_a() 函数；完成后，wave 转到 do_b() 函数，此时，剩余的工作项被屏蔽掉。对于任一功能，只有部分工作项处于事件状态。

回到循环问题，由于循环是分支操作，如果循环条件对某些工作项为真，那么就会出现上述情况，其中一些工作项执行循环中的代码，而其他工作项会被屏蔽掉。但是，在您的代码中:

for(int jb=0; jb < nb; jb++) { /* Foreach block ... */
      pblock[ti] = pos_old[jb*nt+ti]; /* Cache ONE particle position */
      barrier(CLK_LOCAL_MEM_FENCE); /* Wait for others in the work-group */

      for(int j=0; j<nt; j++) { /* For ALL cached particle positions ... */

循环条件不依赖于工作项 ID，这意味着所有工作项将具有完全相同的循环条件，因此它们将遵循相同的执行路径并始终并行运行。

2) In this code snippet, how does the outer and inner loops execute? Does OpenCL know that the outer loop is dividing the work among all the work groups and that the inner loop is trying to divide the work among work-items within each work group?

如对 (1) 的回答所述，由于所有工作项的外循环和内循环的循环条件相同，因此它们始终并行运行。

在 OpenCL 中的工作负载分配方面，完全依赖于开发人员来指定如何分配工作负载。 OpenCL 不知道如何在工作组和工作项之间分配工作量。您可以通过使用全局工作 ID 或本地工作 ID 分配不同的数据和操作来对工作负载进行分区。例如，

unsigned int gid = get_global_id(0);
buf[gid] = input1[gid] + input2[gid];

此代码要求每个工作项从连续内存中获取两个数据并将计算结果存储到连续内存中。

3) If the inner loop is divided among the work-items (meaning that the code within the for loop is executed in parallel, or at least attempted to), how does the addition at the end work? It is essentially doing a = a + f*d, and from my understanding of pipelined processors, this has to be executed sequentially.

     float4 d = p2 - p;
     float invr = rsqrt(d.x*d.x + d.y*d.y + d.z*d.z + eps);
     float f = p2.w*invr*invr*invr;
     a += f*d; /* Accumulate acceleration */

这里，a、f 和 d 在内核代码中定义，没有说明符，这意味着它们仅对工作项本身是私有(private)的。在 GPU 中，这些变量将首先分配给寄存器；然而，寄存器在 GPU 上通常是非常有限的资源，所以当寄存器用完时，这些变量将被放入私有(private)内存中，这称为寄存器溢出(取决于硬件，它可能以不同的方式实现；例如，在某些平台，私有(private)内存是使用全局内存实现的，因此任何寄存器溢出都会导致性能大幅下降)。

由于这些变量是私有(private)的，所以所有的工作项仍然并行运行，并且每个工作项都维护和更新自己的 a、f 和 d，而不会相互干扰。