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)和只读缓存。 (即__ldg())在cc3.5和更高的GPU。使用这些缓存可能在以下两方面有所帮助：

1. 对于一些会破坏线性组织L1的访问模式，您可能会在纹理或只读缓存中获得一些缓存命中，因为这些缓存采用了不同的缓存策略。
2. 在也在使用L1缓存的设备上，通过备用缓存路由“随机”通信将使L1“不受污染”，因此不太可能被破坏。换句话说，L1仍然可以为其他访问提供缓存好处，因为它没有被随机访问破坏。

注意，如果您按照描述的__ldg()使用const __restrict__来修饰适当的指针，编译器可能会通过只读缓存为您路由流量，而不会显式地使用这里。

您还可以在加载和存储上使用缓存控制提示。

类似于上面关于保护L1的建议，在某些设备上它可能是有意义的，在某些情况下，以“非缓存”的方式执行负载和存储。您通常可以通过使用volatile 关键字让编译器为您处理这个问题。您可以将一个普通指针和一个volatile指针保存到相同的数据，以便您可以正则化的访问可以使用“普通”指针，而“随机”访问可以使用volatile版本。实现非缓存访问的其他机制是使用ptxas编译器开关 (例如-Xptxas dlcm=cg)，或者通过适当使用内联PTX和适当的缓存修饰符来管理加载/存储操作。

“非公开”的建议是我可以为“随机”写的主要建议。使用表面机构可能会为某些访问模式提供一些好处，但我认为它不太可能对随机模式作出任何改进。

(*)这种情况在CUDA的最新版本和最近的GPU家庭(如Ampere (cc 8.x) )中发生了变化。有一种新的功能可以将L2的一部分保留给数据持久化。还请参见这里