答案位于答案末尾的 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 毫秒
结果[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。