Ple*_*t94 0 memory pointers memory-management cuda
我不太了解__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__M 上的关键字意味着它是只读的和常量的,那么两者之间的区别是什么,分配的类型？
__restrict__这里nvcc记录了所使用的内容.(请注意,包括gnu编译器在内的各种c ++编译器也支持这个确切的关键字,并且类似地使用它).
它与C99 restrict关键字具有基本相同的语义,C99 关键字是该语言标准的官方部分.
简而言之,__restrict__是程序员与编译器签订的合同,粗略地说,"我只会使用这个指针来引用基础数据".从编译器的角度来看,这取消了表格的一个关键因素是指针别名,这可能会阻止编译器进行各种优化.
如果您想要关于restrictor 的确切定义的更长的正式论文__restrict__,请参考我已经给出的一个链接,或做一些研究.
因此,__restrict__出于优化目的,通常对支持它的编译器很有用.
对于计算能力3.5或更高的设备,这些设备具有称为只读高速缓存的单独高速缓存,其独立于正常L1类型高速缓存.
如果您同时使用__restrict__和const装饰传递给内核的全球指针,那么这也是一个强烈的暗示,编译器,用于cc3.5及更高版本的设备生成代码时,为了使这些全球内存负载通过只读缓存流.这可以提供应用程序性能优势,通常只需很少的其他代码重构.这并不能保证只读缓存的使用,如果它能满足必要的条件,编译器通常会尝试积极地使用只读缓存,即使你不使用这些装饰器也是如此.
__constant__是指GPU上的不同 硬件资源.有很多不同之处:
__constant__ 适用于所有GPU,只能在cc3.5及更高版本上使用只读缓存__constant__标记分配的内存(包含在指定内存分配的行中)限制为最大64KB.只读缓存没有这样的限制.我们没有设置__restrict__分配内存的线路; 它用于装饰指针.__constant__机制OTOH期望所谓的统一访问以获得最快的性能.统一访问本质上意味着warp中的每个线程都在请求来自相同位置/地址/索引的数据.这两种__constant__内存,并标有全局内存const传递给内核代码的指针装饰,是只读的,从内核代码的角度.
无论是使用__restrict__还是其他任何内容,我都没有在您显示的代码中看到任何明显的问题.我唯一的评论是,为了获得最大的好处,您可能想要在内核声明/原型中装饰N和P指针__restrict__,以获得最大的好处,如果这是您的意图.(你不会装饰P用const,很明显.)