为什么这个简单的 OpenCL 代码没有向量化？

Question

下面的代码没有向量化。用'istart = n * 1;'而不是 'istart = n * niters;'确实如此。用'istart = n * 2;'它又没有。

// Kernel for ERIAS_critical_code.py

__kernel void pi(
   int                niters,   
   __global float*    A_d,
   __global float*    S_d,
   __global float*    B_d)                        
{                                                          
   int num_wrk_items  = get_local_size(0);                 
   int local_id       = get_local_id(0); // work item id                  
   int group_id       = get_group_id(0); // work group id                  

   float accum = 0.0f;                              
   int i, istart, iend, n;                                      

   n= group_id * num_wrk_items + local_id;

   istart = n * niters;
   iend   = istart + niters;

   for (i= istart; i< iend; i++){
       accum += A_d[i] * S_d[i];
} 

   B_d[n] = accum;

   barrier(CLK_LOCAL_MEM_FENCE); // test: result is correct without this statement            
}

如果代码不能被矢量化，我会得到：

内核未矢量化

如果可以的话：

内核向量化成功（8）

知道为什么它没有被矢量化吗？

Answer 1

当 niters 为 1 时，它只使 for 循环循环一次。这意味着每个工作项在对内存的合并访问中计算自己的元素。

合并访问是将 N 个相邻线程/工作项映射到 SIMD 硬件（例如宽度为 8）的条件之一。

当 niters 大于 1 时，每个工作项仅适用于相邻工作项之间的 niters 步幅。这意味着 SIMD 硬件毫无用处。每个工作项一次只使用 1 个内存单元。

当 niters 为 2 时，至少只会发生 2 次内存条冲突。但是在非常大的 niters 值下，内存库冲突发生得更多，因此速度非常慢。是否使用 SIMD 无关紧要（是否矢量化），因为它的性能将被锁定在序列化的内存读/写延迟中。

那个 for 循环正在做一个连续的缩减。你应该让它平行。那里有很多例子，选择一个并应用于你的算法。例如，让每个工作项计算 id 和 id+niters/2 之间的总和，然后在 id 和 id+niters/4 上减少它们并像这样继续，直到最后只有 1 个工作项对 id 和 id+1 元素进行最终求和。

如果归约是全局版本，那么您可以对每个工作组进行局部归约，然后将它们的结果以相同的方式应用于另一个内核以进行全局归约。

由于您只对每个工作项进行部分求和，因此您可以执行“每个工作项的跨步求和”，这样每个工作项使用相同的 for 循环但跳跃 M 个元素（其中 M 是某些东西）不会干扰内核工作项上的 SIMD 映射。也许 M 可能是全局元素数（N）的 1/100，for 循环将循环 100 次（或 N/M 次）。像这样的：

                time 1      time 2      time 3      time 4
workitem 1        0           15          30          45
workitem 2        1           16          31          46
workitem 3        2           17          32          47
 ...
workitem 15       14          29          44          59
               coalesced    coalesced   coalesced    coalesced

使用 15 个工作项完成 60 个元素的 15 个部分求和。如果 SIMD 长度可以容纳这 15 个工作项，那就太好了。

最后，不需要屏障操作，因为内核结束是其中所有工作项的全局隐式同步点。仅当您需要在同一内核中的另一个工作项上使用这些书面结果时才需要屏障。