尺寸和对齐要求:
全局存储器指令支持读取或写入大小等于1,2,4,8或16字节的字。 对存在于全局存储器中的数据的任何访问(通过变量或指针)编译为单个全局存储器指令当且仅当数据类型的大小是1,2,4,8或16字节并且数据是自然的 对齐(即,其地址是该尺寸的倍数)。
如果此大小和对齐要求未满足,则访问会使用交织访问模式编译为多条指令,以防止这些指令完全合并。 因此,建议使用满足此要求的类型来存放驻留在全局内存中的数据。
对齐要求自动满足内置类型的char,short,int,long,longlong,float,double如float2或float4。
对于结构,编译器可以使用对齐说明符__align __(8)或__align __(16)来强制执行大小和对齐要求,例如
struct __align__(8) {
float x;
float y;
};
或者:
struct __align__(16) {
float x;
float y;
float z;
};
驻留在全局内存中或由驱动程序或运行时API的内存分配例程之一返回的变量的任何地址始终对齐至少256个字节。
读取非自然对齐的8字节或16字节的字会产生不正确的结果(由几个字组成),因此必须特别注意保持这些类型值的任何值或数组值的起始地址的对齐。 一个典型的例子很容易被忽略,那就是当使用一些自定义的全局内存分配方案时,多个数组的分配(通过多次调用cudaMalloc()或者cuMemAlloc())被分配一个大的内存块 分割成多个数组,在这种情况下,每个数组的起始地址偏离该数据块的起始地址。
二维数组:
一个通用的全局内存访问模式是当索引的每个线程(tx,ty)使用以下地址访问位于类型为type *的地址BaseAddress的宽度二维数组的一个元素时(其中类型满足Maximize利用):
BaseAddress + width * ty + tx
为了使这些访问完全合并,线程块的宽度和阵列的宽度必须是经线大小的倍数。
特别是,这意味着如果一个宽度不是这个大小倍数的数组实际上被分配了一个四舍五入到这个大小的最接近的倍数的宽度,并且相应地填充了它的行,那么这个数组将被更有效地访问。 参考手册中介绍的cudaMallocPitch()和cuMemAllocPitch()函数和相关的存储器复制函数使程序员能够编写非硬件相关的代码来分配符合这些约束的数组。
常量内存:
常量内存空间驻留在设备内存中,并缓存在常量缓存中。
然后将请求分割成多个独立的请求,因为初始请求中有不同的内存地址,吞吐量减少了一个等于单独请求数的因子。
然后在高速缓存命中的情况下以恒定高速缓存的吞吐量来处理所产生的请求,否则以设备存储器的吞吐量来处理所产生的请求。