内存声明：

设备和常量内存：使用__device__或__constant__内存空间说明符在文件范围内声明的内存在使用设备运行时时具有相同的行为。 所有内核都可以读取或写入设备变量，无论内核是由主机还是设备运行时初始启动。 等同地，所有内核将具有与在模块范围内声明的__constant__相同的视图。纹理和表面内存：CUDA支持动态创建的纹理和表面对象1，其中可以在主机上创建纹理参考，传递给内核，由内核使用，然后从主机中销毁。 设备运行时不允许从设备代码中创建或销毁纹理或表面对象，但可以在设备上自由使用并自由传送由主机创建的纹理和表面对象。 无论它们在哪里创建，动态创建的纹理对象总是有效的，并且可以从父级传递给子级内核。设备运行时不支持从设备启动的内核中的传统模块范围（即费米架构）纹理和表面。 模块范围（遗留）纹理可以从主机创建并在设备代码中用于任何内核，但只能由顶级内核（即从主机启动的）使用。

1：动态创建的纹理和曲面对象是CUDA 5.0引入的CUDA内存模型的补充。 有关详细信息，请参阅CUDA编程指南。

共享内存变量声明：在CUDA中，C / C ++共享内存可以声明为静态大小的文件范围变量或函数范围变量，也可以声明为外部变量，其大小由内核的调用者在运行时通过启动配置参数确定。 这两种类型的声明在设备运行时间下均有效。

__global__ void permute(int n, int *data) {    extern __shared__ int smem[];    if (n <= 1)        return;    smem[threadIdx.x] = data[threadIdx.x];    __syncthreads();    permute_data(smem, n);    __syncthreads();    // Write back to GMEM since we can't pass SMEM to children.    data[threadIdx.x] = smem[threadIdx.x];    __syncthreads();    if (threadIdx.x == 0) {        permute << < 1, 256, n / 2 * sizeof(int) >> >(n / 2, data);        permute << < 1, 256, n / 2 * sizeof(int) >> >(n / 2, data + n / 2);    }}void host_launch(int *data) {    permute << < 1, 256, 256 * sizeof(int) >> >(256, data);}

符号地址：

由于所有全局范围的设备变量都在内核的可见地址空间中，所以设备端符号（即标记为__ device_的那些符号）可以通过简单的＆运算符从内核中引用。 这也适用于__constant__符号，但在这种情况下，指针将引用只读数据。考虑到设备端符号可以直接引用，引用符号的CUDA运行时API（例如cudaMemcpyToSymbol（）或cudaGetSymbolAddress（））是多余的，因此设备运行时不支持。 注意这意味着即使在子内核启动之前，也不能在运行的内核中更改常量数据，因为对__constant__空间的引用是只读的。