__restrict__标签的用法

Question

我不太理解数据自动化系统中__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

__restrict__是由nvcc使用的这里文档。(请注意，各种c++编译器，包括gnu编译器也支持这个确切的关键字，并以类似的方式使用它)。

它的语义与C99 restrict关键字(即该语言标准的官方部分 )的语义基本相同。

简单地说，__restrict__是程序员与编译器达成的契约，它说，大概是，“我将只使用这个指针来引用底层数据”。从编译器的角度来看，这从表中取出的一个关键因素是指针混叠，它可以阻止编译器进行各种优化。

如果您想要一篇关于restrict或__restrict__的详细定义的较长的正式论文，请参考我已经提供的链接之一，或者做一些研究。

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

对于计算能力3.5或更高的设备，这些设备有一个称为只读缓存的单独缓存，它独立于普通的L1类型缓存。

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

__constant__指的是一个不同的 GPU上的硬件资源。有许多不同：

• __constant__在所有GPU上都可用，只读缓存仅在cc3.5及更高版本上可用。
• 使用__constant__标记分配的内存(包括在指定内存分配的行中)被限制在最大64 to。只读缓存没有这样的限制。我们不将__restrict__放在分配内存的行上；它用于修饰指针。
• 缓存在只读缓存中的数据具有典型的全局内存访问注意事项--通常我们希望相邻的和连续的访问，以便通过只读缓存最好地合并全局内存读取。__constant__机制OTOH期望以最快的性能实现所谓的统一访问.统一访问本质上意味着，一个翘曲中的每个线程都请求来自同一个位置/地址/索引的数据。

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

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