工作项缓存的全局内存

Question

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

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

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

（我的目标是 AMD GPU）

我需要保证没有的最小尺寸是多少工作项之间有任何并发问题吗？

由于 AMD 有 <= 64 个 CU，我的猜测是

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

Answer 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

然后改内核为interleaved类型（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

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

工作项缓存的全局内存

Global memory for work item cache

opencl