pointers - CUDA:__restrict__ 标签使用

Question

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

我用 __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);

}

我在使用 __restrict__以正确的方式标记N，即只读？
此外，我读过关键字 __constant__ on M 的意思是只读和常量，那么它们两者有什么区别，分配的类型呢？

Answer 1

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

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

简而言之，__restrict__是您作为程序员与编译器签订的契约(Contract)，上面写着 大致 , “我只会使用这个指针来引用底层数据”。从编译器的角度来看，这样做的关键之一是指针别名，它可以阻止编译器进行各种优化。

如果您想要关于 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__，我都没有在您显示的代码中看到任何明显的问题。或其他任何东西。我唯一的意见是，为了获得最大利益，您可能希望同时装饰 N和 P在你的内核声明/原型(prototype)中使用 __restrict__ 的指针，为了最大的利益，如果这是你的意图。 (很明显，你不会用 P 装饰 const。)