CUDA：__restrict__ 标签使用

Question

我不太明白CUDA中__restrict__标签的概念。

我读到使用__restrict__ 可以避免指针别名，特别是，如果指向的变量是只读的，则变量的读取会得到优化，因为它是缓存的。

这是代码的简化版本：

__constant__ float M[M_DIM1][M_DIM2];

__host__ void function(float N[][DIM2], float h_M[][M_DIM2], float P[][DIM2]);

__global__ void kernel_function(const float* __restrict__ N, float *P);

__host__ void function(float N[][DIM2], float h_M[][M_DIM2], float P[][DIM2]) {

    int IOSize = DIM1 * DIM2 * sizeof(float);
    int ConstSize = M_DIM1* M_DIM2* sizeof(float);
    float* dN, *dP;
    cudaMalloc((void**)&dN, IOSize);
    cudaMemcpy(dN, N, IOSize, cudaMemcpyHostToDevice);

    cudaMemcpyToSymbol(M, h_M, ConstSize);

    cudaMalloc((void**)&dP, IOSize);

    dim3 dimBlock(DIM1, DIM2);
    dim3 dimGrid(1, 1);

    kernel_function << <dimGrid, dimBlock >> >(dN, dP);

    cudaMemcpy(P, dP, IOSize, cudaMemcpyDeviceToHost);

    cudaFree(dN);
    cudaFree(dP);

}

我是否以正确的方式在 N 上使用__restrict__ 标签，即只读标签？ 另外，我读到M上的关键字__constant__的意思是只读和常量，那么这两者有什么区别，分配的类型呢？

Answer 1

nvcc 使用的__restrict__ 记录在here 中。 （请注意，包括 gnu 编译器在内的各种 c++ 编译器也支持这个确切的关​​键字，并且类似地使用它。

它与 C99 restrict 关键字的语义基本相同，即an official part of that language standard。

简而言之，__restrict__ 是您作为程序员与编译器签订的合同，大致说，“我只会使用此指针来引用基础数据”。从编译器的角度来看，这样做的关键之一是指针别名，它可以阻止编译器进行各种优化。

如果您想要关于restrict 或__restrict__ 的确切定义的更长的正式论文，请参考我已经提供的链接之一，或进行一些研究。

因此，__restrict__ 通常对支持它的编译器很有用，用于优化目的。

对于计算能力 3.5 或更高版本的设备，这些设备有一个单独的缓存，称为 read only cache，它独立于正常的 L1 类型缓存。

如果您同时使用__restrict__ 和const 来装饰传递给内核的全局指针，那么这也是对编译器的强烈提示，在为cc3.5 和更高版本的设备生成代码时，会导致那些全局内存负载流过只读缓存。这可以提供应用程序性能优势，通常几乎不需要其他代码重构。这并不能保证只读缓存的使用，如果它可以满足必要的条件，编译器通常会尝试积极使用只读缓存，即使你不使用这些装饰器.

__constant__ 指的是一个不同的 hardware resource on the GPU。有很多不同：

• __constant__ 适用于所有 GPU，只读缓存仅适用于 cc3.5 及更高版本
• 使用__constant__ 标记（包含在指定内存分配的行中）分配的内存限制为最大64KB。只读缓存没有这样的限制。我们不会将__restrict__ 放在分配内存的行上；用于装饰指针。
• 缓存在只读缓存中的数据具有典型的全局内存访问注意事项 - 通常我们需要相邻和连续访问，以便通过只读缓存最佳地合并全局内存读取。 __constant__ 机制 OTOH 期望所谓的统一访问以获得最快的性能。统一访问本质上意味着 warp 中的每个线程都从相同的位置/地址/索引请求数据。

__constant__ 内存和在传递给内核代码的指针上标有 const 装饰器的全局内存从内核代码的角度来看都是只读的。

我在您显示的代码中看不到任何明显的问题，无论是使用__restrict__ 还是其他任何东西。我唯一的意见是，为了获得最大利益，您可能希望使用__restrict__ 装饰内核声明/原型中的N 和P 指针，以获得最大利益，如果这是您的意图。 （很明显，你不会用const 装饰P。）