工作项缓存的全局内存

Posted 2023-02-16

【中文标题】工作项缓存的全局内存

【英文标题】：Global memory for work item cache

【发布时间】：2017-03-09 15:28:49

【问题描述】：

由于本地内存限制，我需要使用全局内存作为工作项的缓存。

假设我有 1000 个工作组，每个工作组有 64 个工作项。每个项目都需要 4K 缓存。工作项完成后，缓存不需要保留。

我将分配一个单独的全局内存缓冲区并为工作项分配 4K 块。

（我的目标是 AMD GPU）

我需要保证不会出现的最小尺寸是多少 工作项之间的任何并发问题？

由于 AMD 有

64 * 128 * 4000 字节，并使用（全局工作项 ID % (64*128)） 将缓存块分配给工作项。

【问题讨论】：

您可以尝试使用 codexl 对其进行分析。它告诉你关于内核的必要条件，瓶颈在哪里

【参考方案1】：

如果每个缓存项（由global work item ID % (64*128) 访问）是一个 4000 字节长的结构，并且如果实现不强制每个结构在 4096 字节上对齐，并且如果缓存行大小不是 4000 的精确除数如果全局内存条的步长不是 4000 的精确除数，那么这应该不是问题。

---

用 codexl 分析这个内核，（16k 个工作项需要 0.5 秒）：

    __kernel void test(__global float * a)
    
        int i=get_global_id(0)*4096;
        for(int j=0;j<4096;j++)
            a[i+j]*=2.0f;
    

还有一些输出：

内存单元停止 %55 缓存命中 %45 内存单元忙 %99 值忙 %0.05

然后将内核改为交错式（0.25s执行）：

    __kernel void test(__global float * a)
    
        int i=get_global_id(0);
        for(int j=0;j<4096;j++)
            a[i+j*4096*4]*=2.0f;
    

内存单元停止 %57 缓存命中 %47 内存单元忙 %84 值忙 %1.5

因此交错模式对内存单元的压力较小，命中缓存的频率更高，而 ALU 部件的馈送频率更高，完成速度更快 %50。

然后试试这个：

__kernel void test(__global float * a)

    int i=get_global_id(0)*4100;
    for(int j=0;j<4100;j++)
        a[i+j]*=2.0f;

这花费了 0.37 秒，比 4096 版本快 %30，但内存单元停顿率更高（端点不对齐一定导致这在不必要的数据获取上浪费了一些周期）并且缓存命中减少到 %37。

测试 GPU 为 R7-240

---

最后一次使用结构的测试：

typedef struct test_struct

   float test_field[4096];
strr;
__kernel void test(__global strr * a)

    int i=get_global_id(0);
    for(int j=0;j<4096;j++)
    a[i].test_field[j]*=2.0f;

这在 0.53 秒内完成，并且在开始时具有与跨步内核相似的分析数据。

空内核在 0.25 秒内执行，因此它不会加载整个结构。只读取需要的元素。

---

以组为中心的交错全局访问分析：

typedef struct test_struct

   float test_field[4096];
strr;
__kernel void test(__global strr * a)

    int iLocal=get_local_id(0);
    int iGroup=get_group_id(0);
    for(int j=0;j<64;j++)
    a[iGroup].test_field[iLocal+j*64]*=2.0f;

又是 0.25 秒，所以它尽可能快。

缓存命中：%44 内存单元忙：%82 内存单元停止：%67 值忙：%0.9

所以它拥有最好的条件，即使没有缓存。

【讨论】：

谢谢。我的意思是 4096 字节的缓存块。但是，是否可以保证根据全局工作项 id 按特定顺序安排工作项？如果是这种情况，那么我可以确定我没有两个工作项使用相同的缓存块

如果一个工作项可以使用一个计算单元中的所有线路资源，那应该不是问题，但我不太愿意这样做，因此他们更有可能并行访问。我不能对此有任何保证。 128x64 是 8k，因此比 8k 更多的工作项将对同一个缓存元素进行多次访问（大 gpu 可以在多个时钟内拥有它，具有 4k 个内核）

谢谢。那么，您是说为了安全起见，我应该为每个工作项分配一个 4096 块吗？我想为多个工作项重复使用相同的缓存块，但前提是可以保证工作项不会同时访问该块

AMD 自己的论坛说 2 的大功率是不好的。因为他们说只使用了几行。 4000应该是好的或4103也许。克服这样的事情。但是您也可以使用交错缓存。这样第一个工作项访问第 1 个字节，第 10000 个字节 ..... 第 4096000 个字节，用于其自身所有 4096 字节的缓存数据。这将允许显式内存请求平衡，如果不是隐式的话。

非常感谢您分析这些内核。对不起，我想我可能没有清楚地解释这个问题。每个工作组有 64 个工作项，所有工作项共享一个缓存缓冲区。访问被设计为跨工作项合并读取。即每个工作项在缓存中拥有一条垂直的数组点行。所以，项目 0 拥有 0,64,128，......，项目 1 拥有 1,65,129，等等。所以，读取被合并。