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代码产生负面的性能影响。