由于本地内存限制，我需要使用全局内存作为工作项的缓存。
假设我有 1000 个工作组，每个工作组有 64 个工作项。每个项目都需要 4K 缓存。工作项完成后，缓存不需要保留。
我将分配一个单独的全局内存缓冲区并为工作项分配 4K 块。
（我的目标是 AMD GPU）
我需要保证不会出现的最小尺寸是多少 工作项之间的任何并发问题？
由于 AMD 有
64 * 128 * 4000 字节，并使用（全局工作项 ID % (64*128)） 将缓存块分配给工作项。
【问题讨论】：
标签： opencl
如果每个缓存项（由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;
}
还有一些输出：
然后将内核改为交错式（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;
}
因此交错模式对内存单元的压力较小，命中缓存的频率更高，而 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
所以它拥有最好的条件，即使没有缓存。
【讨论】：