带有 OpenCL 的 Intel HD 6000 本地内存带宽 [关闭]

Posted 2023-02-17

【中文标题】带有 OpenCL 的 Intel HD 6000 本地内存带宽 [关闭]

【英文标题】：Intel HD 6000 local memory bandwidth with OpenCL [closed]

【发布时间】：2016-09-06 21:24:13

【问题描述】：

我正在 OpenCL 中进行一些本地/全局内存优化；在查看两年前的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 浮点值会导致读取 64G 内存；内核的执行耗时 453945 μs，估计本地内存带宽为 151 GB/s。

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

更令人困惑的是，我在一些专用的中档 GPU 上的性能越来越差：AMD R9 m370x 和 Nvidia GT 750m 都需要 700-800 毫秒。这些卡比英特尔的 HD 6000 稍旧，因此可能与它有关。

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

【问题讨论】：

我认为 3-4TB/s 的本地内存太疯狂了。我不知道如果来自哪里，但该设备只能达到 700 GFLOP/s，实际上不可能有 4TB/s 的传输速度。从硬件设计的角度来看，这意味着内存被过度设计，能够提供比 ALU 处理的数据多 4 倍的数据。 4Tb/s 或 200GB/s 会更有意义。与您获得的类似。

接近带宽限制，无法完成一次计算；并且在接近计算限制时类似，无法完成单个缓存操​​作。

【参考方案1】：

Intel HD 6000 有两个切片，每个切片包含三个子切片，每个子切片分别连接到共享本地内存（参见此处的图表 http://www.notebookcheck.net/Intel-HD-Graphics-6000.125588.0.html），每个周期的带宽为 64 字节，因此假设 1 GHz 时钟，您将获得 6 * 64 * 1 GHz = 384 GB/s 来自本地内存的峰值带宽。如果您访问 16 个本地内存库中的每一个（本地内存是高度存储的，因此您可以在每个周期独立地从每个库中获取 4 个字节），您就会明白这一点。您可以通过 loc[id] 访问或类似的方式获得这种模式。下载面向 OpenCL 的英特尔 SDK https://software.intel.com/en-us/intel-opencl - 它为您提供了程序集视图：您的代码将编译为 SIMD32，但您的代码生成的程序集非常糟糕，因为您不断从每个 SIMD 通道敲击相同的位置，所以您很幸运，您获得了高达 151 GB/s 的速度。

【讨论】：

谢谢，这是很好的信息吗？我将如何从没有列出“切片”和“子切片”等架构数据的 Nvidia 卡（如 GTX 1070）计算预期（峰值）本地内存带宽？

对不起，Lukas，我是 Intel 人 :) 也许这里的其他人知道 Nvidia 硬件。

【参考方案2】：

答案位于答案末尾的 edit2 部分。

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

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

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

loc[23] = 34;

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

还有可能

for(int i = 0; i

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

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

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

---

编辑：以下代码在 AMD-R7-240 卡上以 900MHz（不超过 30 GB/s 内存和 600 GFlops）运行，得到 1600 万个结果元素。

        __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 毫秒计算 +（映射/取消映射而不是读/写） 240 毫秒的计算时间（264 GB/s 本地内存带宽）

Intel(R) HD Graphics 400

@ 600 MHz (45 GB/s) 为 1410 毫秒 警告：这是__global数组访问

结果[id] = ...

__全局数组访问是此算法的此设备的瓶颈。

230 毫秒而不是 HD 400 的 1410 毫秒 !!!!! （这应该是缓存/本地带宽）

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。

【讨论】：

感谢您的彻底（和快速）回复！您提到使用 __private 寄存器，如果我理解正确，在这种情况下将是 rtemp，iGPU 不支持并强制 GPU-CPU 通信——这将如何解释 nvidia GPU 上的类似性能？感谢您的 float8 建议；它似乎不会影响性能，但我会尝试一些不同的访问方法，看看会发生什么。省略“loc[23] = 34;”似乎也不影响性能；执行所需的时间与本地数组的默认值为零相同。

我将 rtemp 移动到本地内存，执行时间加倍；将其改回私有削减执行时间，这有望排除私有寄存器成为瓶颈。

在检查了 glGetDeviceIDs 和 clGetDeviceInfo 之后，我在我测试的机器上使用了集成的英特尔 GPU，而不是 Nvidia / AMD dGPU，这导致时间比我的快 35% 左右——在这个有一个强有力的理由表明 150 GB/s 是典型的本地内存速度。您使用的图形是您见过的，还是只是一个示例来说明您的观点？

@LukasPalmer 在最后添加了答案，请查收。

太棒了！在我的情况下，即使没有优化，全局内存似乎仍然只有本地内存的一半左右，但这是一些很好的信息！我没想到所有线程都在同一个位置，idL 技巧解决了这个问题。