优化CUDA中的随机访问读取和随机访问写入

Question

我想在以下代码中优化随机访问读取和随机访问写入:

__global__ void kernel(float* input, float* output, float* table, size_t size)
{
int x_id = blockIdx.x * blockDim.x + threadIdx.x;
if (x_id > size)
    return;

float in_f = input[x_id];
int in_i = (int)(floor(in_f));
int table_index = (int)((in_f - float(in_i)) * 1024000.0f );
float* t = table + table_index;
output[table_index] = t[0] * in_f;

}

如您所见,表和输出的索引在运行时确定,并且完全随机.

我知道我可以使用纹理内存或__ldg()读取这些数据.所以,我的问题是:

1. 有没有比使用纹理内存更好的方法来读取随机索引的数据__ldg()？
2. output[table_index]如上所述,随机访问写入怎么样？

实际上,我在这里添加代码以给出随机访问读写的示例.我不需要代码优化,我只需要高级描述处理这种情况的最佳方法.

Answer 1

在GPU上随机访问数据没有神奇的子弹.

最好的建议是尝试执行数据重组或其他一些方法来规范数据访问.对于重复/重度访问模式,即使对数据进行排序操作等密集方法也可能导致性能的整体净改进.

由于您的问题暗示随机访问是不可避免的,因此您可以做的主要是智能地使用缓存.

L2是设备范围的缓存,所有DRAM访问都通过它.因此,如果你有大规模的随机访问,那么L2的颠簸可能是不可避免的.对于读或写访问,没有任何功能可以禁用(有选择地或以其他方式)L2.

对于较小规模的情况,您可以做的主要是通过一个"非L1"缓存路由访问,即纹理缓存(在所有GPU上)和cc3.5及更高版本上的只读缓存(即__ldg())图形处理器.使用这些缓存可能有两种方式:

1. 对于一些会破坏线性组织L1的访问模式,由于这些缓存采用了不同的缓存策略,您可能会在纹理或只读缓存中获得一些缓存命中.

2. 在使用L1缓存的设备上,通过备用缓存路由"随机"流量将使L1"未受污染",因此不太可能发生冲突.换句话说,L1可能仍然为其他访问提供缓存益处,因为它不会被随机访问所淹没.

请注意__ldg(),如果您使用此处const __restrict__所述装饰适当的指针,编译器可能会为您通过只读缓存路由流量,而无需明确使用.

与上述保护L1的建议类似,在某些设备上可能有意义,在某些情况下以"未缓存"的方式执行加载和存储.您通常可以通过使用volatile 关键字让编译器为您处理此问题.你可以保持普通和volatile指向同一数据的指针,这样你可以规范的访问可以使用"普通"指针,"随机"访问可以使用该volatile版本.追求未缓存访问的其他机制是使用ptxas编译器开关(例如-Xptxas dlcm=cg),或者通过适当使用具有适当缓存修饰符的内联PTX来管理加载/存储操作.

"未缓存"的建议是我可以为"随机"写入提供的主要建议.使用表面机制可能会为某些访问模式带来一些好处,但我认为不可能对随机模式进行任何改进.