CUDA学习（三十五）-白红宇

__constant__:

__constant__内存空间说明符（可选）与__device__一起使用，声明一个变量：

驻留在恒定的内存空间中，

是否具有创建CUDA上下文的生命周期，

每个设备都有独特的对象，

可以从网格中的所有线程和主机通过

可以从网格中的所有线程和主机通过运行时库（cudaGetSymbolAddress（）/ cudaGetSymbolSize（）/
cudaMemcpyToSymbol（）/ cudaMemcpyFromSymbol（））。

__shared__:

__shared__内存空间说明符，可以与__device__一起使用，

声明一个变量：

驻留在线程块的共享内存空间中，

有块的生命周期，

每块有一个独特的对象，

只能从块中的所有线程访问
将共享内存中的变量声明为外部数组（如

extern __shared__ float shared[];

数组的大小在启动时确定（请参阅执行配置）。所有以这种方式声明的变量，从内存中的相同地址开始，以便数组中变量的布局必须通过偏移量进行显式管理。例如，如果有人想要相当于:

short array0[128];float array1[64];int array2[256];

在动态分配的共享内存中，可以通过以下方式声明和初始化数组：

extern __shared__ float array[];__device__ void func() // __device__ or __global__ function{    short* array0 = (short*)array;    float* array1 = (float*)&array0[128];    int* array2 = (int*)&array1[64];}

请注意，指针需要与它们指向的类型对齐，因此下面的代码不起作用，因为array1未对齐到4个字节。

extern __shared__ float array[];__device__ void func() // __device__ or __global__ function{    short* array0 = (short*)array;    float* array1 = (float*)&array0[127];}

表3列出了内置矢量类型的对齐要求

__managed__:

__managed__内存空间说明符（可选与__device__一起使用）声明一个变量：

可以从设备和主机代码中引用，例如，可以获取其地址，也可以直接从设备或主机功能读取或写入。

有申请的有效期。
有关更多详细信息，请参阅__managed__ Memory Space Specifier

__restrict__:

nvcc通过__restrict__关键字支持限制指针;

在C99中引入了受限制的指针以缓解存在于C型语言中的混叠问题，并且抑制了从代码重新排序到常见子表达式消除的所有类型的优化

这是一个受到别名问题影响的例子，使用受限制的指针可以帮助编译器减少指令的数量：

void foo(const float* a,    const float* b,    float* c){    c[0] = a[0] * b[0];    c[1] = a[0] * b[0];    c[2] = a[0] * b[0] * a[1];    c[3] = a[0] * a[1];    c[4] = a[0] * b[0];    c[5] = b[0];    ...}

在C型语言中，指针a，b和c可能是别名，所以任何通过c写入都可以修改a或b的元素。这意味着为了保证函数的正确性，编译器不能将[0]和b [0]加载到寄存器中，将它们相乘，并将结果存储到c [0]和c [1]中，因为结果会与抽象执行模型，例如，如果[0]与c [0]确实位于相同的位置。所以编译器不能利用公共的子表达式。同样，编译器不能只将c [4]的计算重新排列到c [0]和c [1]的计算邻近位置，因为前面写入c [3]可能会将输入更改为计算c [4]。

通过制作a，b和c限制指针，程序员向编译器断言指针实际上不是别名，在这种情况下意味着通过c写入不会覆盖a或b的元素。这改变了函数原型如下：

void foo(const float* __restrict__ a,    const float* __restrict__ b,    float* __restrict__ c);

请注意，所有指针参数都需要被编译器优化器限制，以获得任何好处。通过添加__restrict__关键字，编译器现在可以重新排序并随意执行常见的子表达式删除，同时保留与抽象执行模型相同的功能：

void foo(const float* __restrict__ a,    const float* __restrict__ b,    float* __restrict__ c){    float t0 = a[0];    float t1 = b[0];    float t2 = t0 * t2;    float t3 = a[1];    c[0] = t2;    c[1] = t2;    c[4] = t2;    c[2] = t2 * t3;    c[3] = t0 * t3;    c[5] = t1;    ...}

这里的效果是减少了内存访问次数，减少了计算次数。这是由于“缓存”的负载和常见的子表达式的登记压力的增加而平衡的。

由于寄存器压力在许多CUDA代码中是一个关键问题，由于占用率降低，使用受限制的指针可能会对CUDA代码产生负面的性能影响。