Intel HD 6000 本地内存带宽与 OpenCL

Question

我正在 OpenCL 中进行一些 local/global 内存优化；在查看了两年前的 this question 之后，我认为我做错了什么，因为本地内存 IO 似乎比应有的速度慢得多。我的 GPU 是 Intel HD 6000。

这是我的测试设置，内核源代码：

__kernel void vecAdd(__global float* results, const unsigned int n, __local float* loc)
{
   int id = get_global_id(0);
   if(id < n) {
      float rtemp = 0;
      loc[23] = 34;
      for(int i = 0; i < 1024; i ++) {
         rtemp += loc[(i * 445) % 1024];
      }
      results[id] = rtemp;
   }
}

内核所做的就是获取本地浮点数组 loc 并将其中的随机值添加到全局输出向量中。片段“(i * 445) % 1024”用于保证随机访问本地内存；性能比最后提到的没有随机化的数字好一点（~30% 加速）。

我将内核排队等待 16777216 / 16M 次迭代，工作组大小为 256，本地缓冲区为 1024 个浮点数，除 l[23] 外全为零。
总的来说，这使得总共16M * 1 = 16M写入和16M * 1024 = 16G读取到本地内存。
还有大约 16M * 1024 * 2 个浮点运算，可能更多取决于模数的计算方式，但 HD 6000 的浮点性能约为 768 GFLOPS，这不应该是瓶颈.
16G读取float值导致读取64G内存；内核的执行需要 453945 微秒才能完成，估计本地内存带宽为 151 GB/s.

引用问题中的数字表明，现代显卡（自 2014 年起）的内存带宽可能比我在我的机器上测得的要高得多；文章中引用的数字（可能是为了比较而随机示例）是 3-4 TB/s；虽然我的卡是集成卡而不是专用卡，但考虑到它在 2015 年发布，这似乎仍然是一个缓慢的数字。

让事情变得更加混乱的是，我在一些专用的中端 GPU 上的性能越来越差：AMD R9 m370x 和 Nvidia GT 750m 都花费了 700-800 毫秒。这些卡比 Intel 的 HD 6000 稍微老一些，所以这可能与它有关。

是否有任何潜在的方法可以从本地内存中榨取更多性能，或者我是否尽可能有效地利用本地内存？

Answer 1

答案在答案结尾的 edit2 部分。

如果专用 gpu 时序不佳，您可以尝试流水线读取+计算+写入操作，例如

从左到右，它在第二步开始重叠操作，因此隐藏了计算延迟，然后第三步也隐藏了写入延迟。这是将可分离的作品分成 4 部分的示例。也许更多的部分会产生更慢的结果，应该对每个设备进行基准测试。内核执行只是一个 "adding"，所以它总是隐藏的，但更重的可能不是。如果该显卡可以同时进行读取和写入，这会减少 I/O 延迟。图片还显示了空闲（垂直为空）的时间线，因为冗余同步使其比打包但更快的版本更具可读性。

您的 igpu 151 GB/s 带宽可能是 cpu-缓存。它没有可寻址寄存器 space，因此即使使用 __private 寄存器也可以使其从缓存中获取。每个 cpu 或 gpu.

缓存也有不同的线宽

loc[23] = 34;

有多个线程的竞争条件并被序列化。

并且有可能

for(int i = 0; i < 1024; i ++) { rtemp += loc[(i * 445) % 1024]; }

正在自动展开并对指令缓存和 cache/memory 施加压力。您可以尝试不同级别的展开。

您确定该 igpu 的每个执行单元使用 8 个内核吗？也许每个 EU 仅使用 1 个核心，这可能不足以完全强调 cache/memory（例如通过使用所有第一个核心但仅此而已的缓存行冲突）？尝试使用 float8 版本而不仅仅是 float。最新的英特尔 cpu 每秒超过 1 TB。

GFLOPS 限制很少被接近。优化代码约为 %50，代码不可读约为 %75，无意义代码约为 %90。

编辑： 下面的代码是运行在 900MHz 的 AMD-R7-240 卡上（不超过 30 GB/s 内存和 600 GFlops）对于结果的 16M 个元素。

        __kernel void vecAdd(__global float* results )
        {
           int id = get_global_id(0);
           __local float loc[1024]; // some devices may slow with this
           if(id < (4096*4096)) {
              float rtemp = 0;
              loc[23] = 34;
              for(int i = 0; i < 1024; i ++) {
                 rtemp += loc[(i * 445) % 1024];
              }
              results[id] = rtemp;
           }
        }

花了

575 毫秒（无管道）写入+计算+读取
530 毫秒（2 部分流水线）写入+计算+读取
510 毫秒（8 部分流水线）写入+计算+读取
455 毫秒计算（140 GB/s 本地内存带宽）

Edit2： 优化缓存行利用率、计算简化和着色器核心中更少的气泡：

        __kernel void vecAdd(__global float* results )
        {
           int id = get_global_id(0);
           int idL = get_local_id(0);
           __local float loc[1024];
           float rtemp = 0;
           if(id < (4096*4096)) {

              loc[23] = 34;
           }

           barrier (CLK_LOCAL_MEM_FENCE);

           if(id < (4096*4096)) {
              for(int i = 0; i < 1024; i ++) {
                 rtemp += loc[(i * 445+ idL) & 1023];
              }
              results[id] = rtemp;
           }
        }

325 毫秒（16 部分流水线）写入+计算+读取
270 毫秒计算（235 GB/s 本地内存带宽）

loc[(i * 445) % 1024];

对于所有线程都是相同的，都是随机的，但在每一步都更改为相同的值，通过相同的缓存行访问。向所有线程添加局部变体但最终具有相同的总和，使用更多行。

% 1024

使用

进行了优化

&1023

最后，在 loc[23] = 34 之后消除 SIMD 中任何指令气泡的屏障；

Edit3: 添加一些循环展开并将本地工作组大小从 64 增加到 256（edit 和 edit2 为 64）

        __kernel void vecAdd(__global float* results )
        {
           int id = get_global_id(0);
           int idL = get_local_id(0);
           __local float loc[1024];
           float rtemp = 0;
           float rtemp2 = 0;
           float rtemp3 = 0;
           float rtemp4 = 0;
           if(id < (4096*4096)) {

              loc[23] = 34;
           }

           barrier (CLK_LOCAL_MEM_FENCE);

           if(id < (4096*4096)) {
              int higherLimitOfI=1024*445+idL;
              int lowerLimitOfI=idL;
              int stepSize=445*4;
              for(int i = lowerLimitOfI; i < higherLimitOfI; i+=stepSize) {
                 rtemp += loc[i & 1023];
                 rtemp2 += loc[(i+445) & 1023];
                 rtemp3 += loc[(i+445*2) & 1023];
                 rtemp4 += loc[(i+445*3) & 1023];
              }
              results[id] = rtemp+rtemp2+rtemp3+rtemp4;
           }
        }

290 毫秒（8 部分流水线）写入+计算+读取，无需冗余同步（在其他基准测试中忘记了这一点）
pci-e 2.0 8x 而不是 4x 上为 278 毫秒
4 个队列 (rcw + rcw + rcw + rcw) 没有事件而不是 3 个队列 (r+c+w) 有事件流的 249 毫秒。（每个队列 32 个部分，因此总共 128x rcw 部分）
243 毫秒计算 +（map/unmap 而不是 read/write）
240 毫秒计算（264 GB/s 本地内存带宽）
1410 毫秒 Intel(R) 高清显卡 400 @ 600 MHz (45 GB/s)
- 警告：这是__global数组访问
results[id] = ...

__global array access is bottleneck for this device for this algorithm.

230 ms instead of 1410 ms for HD 400 !!!! (this should be cache/local bandwidth)
12 个计算单元，每个有 8 个内核 =>96 个内核 45 GB/s 表示 1 个内核 0.5 GB/s @600 MHz 或 **每个时钟每个内核近 1 个字节* *
您的 igpu 每 3 个周期每个内核可以读取 1B，但它总共有 384 个内核 => **192 GB/s（您接近限制）**
看这张图，它每片写入 64B，这意味着每周期每 192 个内核读取 64 字节，或者每 3 个周期每 192 个内核读取 192 字节：

- 根据分析器，VGPR 使用将内核占用率限制为 %60。

Answer 2

Intel HD 6000 有两个片，每个片有三个子片，每个子片分别连接到共享本地内存（参见此处的图表 http://www.notebookcheck.net/Intel-HD-Graphics-6000.125588.0.html ) with a bandwidth of 64 bytes per cycle, so assuming 1 GHz clock, you get 6 * 64 * 1 GHz = 384 GB/s of peak BW from local memory. You get that if you are hitting every one of 16 banks of local memory (local memory is highly banked so you could fetch 4 bytes per cycle from each bank independently). You get that kind of pattern with loc[id] access or something like that. Download Intel SDK for OpenCL https://software.intel.com/en-us/intel-opencl - 它为您提供汇编视图等：您的代码将编译为 SIMD32 ，但是您的代码生成的程序集非常糟糕，因为您从每个 SIMD 通道不断敲打相同的位置，所以您很幸运，您得到了高达 151 GB/s.

Intel HD 6000 本地内存带宽与 OpenCL

Intel HD 6000 local memory bandwidth with OpenCL

c

c++

memory-management

gpgpu

opencl