在 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;

}

如您所见，table 和输出的索引在运行时确定，并且完全随机。

我知道我可以使用纹理内存或 __ldg() 来读取此类数据。所以，我的问题是：

有没有比使用纹理内存或 __ldg() 更好的读取随机索引数据的方法？
上面output[table_index]的随机访问写入呢？

其实我这里加的代码是为了给一个随机访问读写的例子。我不需要代码优化，我只需要对处理这种情况的最佳方法进行高级描述。

Answer 1

在 GPU 上随机访问数据没有灵丹妙药。

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

既然你的问题暗示了随机访问是不可避免的，你能做的主要是智能地利用缓存。

二级缓存是设备范围的缓存，所有的DRAM访问都经过它。因此，如果你有大规模的随机访问，L2 的颠簸可能是不可避免的。没有任何功能可以（有选择地或以其他方式）禁用 L2 以进行读取或写入访问 (*)。

对于较小规模的情况，您可以做的主要事情是通过“非 L1”缓存之一路由访问，即 cc3.5 上的 texture cache (on all GPUs) and the Read-Only cache（即 __ldg()）和更高的 GPU。使用这些缓存可能在两个方面有所帮助：

对于某些会破坏线性组织的 L1 的访问模式，您可能会在纹理或只读缓存中获得一些缓存命中，因为这些缓存采用了不同的缓存策略。
在同时使用 L1 缓存的设备上，通过备用缓存路由“随机”流量将使 L1 保持“未受污染”，因此不太可能发生抖动。换句话说，L1 可能仍会为其他访问提供缓存优势，因为它不会被随机访问扰乱。

请注意，如果您按照 here[所述，用 const __restrict__ 装饰适当的指针，编译器可能会为您路由流量通过只读缓存，而无需显式使用 __ldg() =23=]

您还可以在加载和存储时使用。

类似于上述保护 L1 的建议，在某些设备上可能有意义，在某些情况下以“未缓存”方式执行加载和存储。您通常可以通过使用 volatile keyword. You can keep both an ordinary and a volatile pointer to the same data, so that accesses that you can regularize can use the "ordinary" pointer, and "random" accesses can use the volatile version. Other mechanisms to pursue uncached access would be to use the ptxas compiler switches (e.g. -Xptxas dlcm=cg) or else manage the load/store operations via appropriate use of inline PTX with appropriate caching modifiers.

让编译器为您处理这个问题

“未缓存”建议是我可以为“随机”写入提供的主要建议。使用 surface mechanism 可能会为某些访问模式提供一些好处，但我认为它不太可能对随机模式做出任何改进。

(*) 这在最新版本的 CUDA 和最新的 GPU 系列（例如 Ampere (cc 8.x) 中已更改。有一项新功能可以为 data persistence. Also see here

保留一部分 L2

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

Optimizing random access read and random access write in CUDA

cuda

gpgpu