用户
 找回密码
 立即注册
林俊熙 该用户已被删除
发表于 2010-2-9 10:25:12
69502
4.2.2.1  _device_
_device_ 限定符声明位于设备上的变量。
在接下来的三节中介绍的其他类型限定符中,最多只能有一种可与 _device_ 限定符一起使用,以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符,则变量具有以下特征:
        位于全局存储器空间中;
        与应用程序具有相同的生命周期;
        可通过网格内的所有线程访问,也可通过运行时库从主机访问。
4.2.2.2  _constant_
_constant_ 限定符可选择与 _device_ 限定符一起使用,所声明的变量具有以下特征:
        位于固定存储器空间中;
        与应用程序具有相同的生命周期;
        可通过网格内的所有线程访问,也可通过运行时库从主机访问。
4.2.2.3  _shared_
_shared_ 限定符可选择与 _device_ 限定符一起使用,所声明的变量具有以下特征:
        位于线程块的共享存储器空间中;
        与块具有相同的生命周期;
        尽可通过块内的所有线程访问。
只有在 _syncthreads()_(参见第 4.4.2 节)的执行写入之后,才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量,否则只要之前的语句完成,编译器即可随意优化共享存储器的读写操作。
将共享存储器中的变量声明为外部数组时,例如:
extern __shared__ float shared[];
数组的大小将在启动时确定(参见第 4.2.3 节)。所有变量均以这种形式声明,在存储器中的同一地址开始,因此数组中的变量布局必须通过偏移显式管理。例如,如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容:
short array0[128];
float array1[64];
int array2[256];
则应通过以下方法声明和初始化数组:
extern __shared__ char array[];
__device__ void func() // __device__ or __global__ function
{
    short* array0 = (short*)array;
     float* array1 = (float*)&array0[128];
     int* array2 = (int*)&array1[64];
}
4.2.2.4  限制
不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。
_shared_ 和 _constant_ 变量具有隐含的静态存储。
_device_、_shared_ 和 _constant_ 变量无法使用 extern 关键字定义为外部变量。
_device_ 和 _constant_ 变量仅允许在文件作用域内使用。
不可为设备或从设备指派 _constant_ 变量,仅可通过主机运行时函数从主机指派(参见第 4.5.2.3 节和第 4.5.3.6 节)。
_shared_ 变量的声明中不可包含初始化。
在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下,编译器可能选择将其置于本地存储器中。如果使用占用了过多寄存器空间的大型结构或数组,或者编译器无法确定其是否使用固定数量索引的数组,则往往会出现这种情况。检查 ptx 汇编代码(通过使用 –ptx 或 –keep 选项编译获得)即可在初次编译过程中确定一个变量是否位于本地存储器中,因为它将使用 .local 助记符声明,可使用 ld.local 和 st.local 助记符访问。如果不是这样,在后续编译阶段仍能确定是否占用了目标架构的过多寄存器空间。可通过使用 --ptxas- options =-v 选项编译来进行检查,这将报告本地存储器的使用情况(lmem)。
只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间,此类指针即受支持,否则将仅限于指向在全局存储器空间中分配或声明的存储器。
如果取消在主机上执行的代码中全局或共享存储器指针,或者在设备上执行的代码中主机存储器指针的引用,将导致不确定的行为,往往会出现分区错误和应用程序终止。
通过获取 _device_、_shared_ 或 _constant_ 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress() (参见第 4.5.23 节)获取的 _device_ 或 _constant_ 变量的地址仅可在主机代码中使用。
使用道具 举报 回复
发表于 2010-12-23 20:18:48
:lovely:
使用道具 举报 回复 支持 反对
发表于 2010-12-23 20:18:59
:iproud:
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册