优化 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”缓存之一路由访问，即texture cache（在所有 GPU 上）和the Read-Only cache（即__ldg()） cc3.5 及更高版本的 GPU。使用这些缓存可能有两种帮助：

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

2. 在同时使用 L1 缓存的设备上，通过备用缓存路由“随机”流量将使 L1 保持“不受污染”，因此不太可能发生抖动。换句话说，L1 仍然可以为其他访问提供缓存优势，因为它不会受到随机访问的影响。

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

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

与上述保护 L1 的建议类似，在某些设备上可能有意义，在某些情况下以“未缓存”方式执行加载和存储。您通常可以通过使用 volatile keyword 让编译器为您处理此问题。您可以将普通指针和volatile 指针同时保留到相同的数据，以便您可以正则化的访问可以使用“普通”指针，而“随机”访问可以使用volatile 版本。追求非缓存访问的其他机制是使用ptxas compiler switches（例如-Xptxas dlcm=cg）或通过适当使用inline PTX和适当的caching modifiers来管理加载/存储操作。

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

(*) 这在最新版本的 CUDA 和最近的 GPU 系列（例如 Ampere (cc 8.x)）中发生了变化。有一项新功能可以为data persistence 保留一部分 L2。另见here