在合并访问这里，不要跟shared memory的bank conflict搞混淆了，这里很重要。

         首先，设备device有能力，在一个单一指令下，从global memory中读取32-bit, 64-bit, 128-bit字进入寄存器register。

         分配如下：

__device__ type device[32];

type data = device[tid];

编译一个单一加载指令，type必须是sizeof(type) = 4、8 或者 16 这样的，

                                            而且要求内置类型像float2 或 float4 一样的自动完成的。

对于结构、大小和队列要求可以通过编译器强制使用队列指定的__align__(8) 或 __align__(16)，

例如

struct __align(8)__{

float a;

float b;

};

//或者是

struct __align(16)__{
float a;
float b;

float c;
};

对于结构 > 16字节的，编译器生成几个加载指令，来保证它生成最低数量的指令，这样的结构应该用__align__(16)定义

例如：

struct __align(16)__{
float a;
float b;

float c

float d

float e;

};

//!被编译成两个128-bit的加载指令，而不是5个32-bit加载指令

其次，全局内存地址同时被每线程的一个half-warp访问（执行读和写指令）时，应该排列好，以便内存的存取可以结合进入一个接近单一的，排列好的内存存取。

         它意味着，在每一个half-warp中，在half-warp中的第N个线程应该访问该地址。

HalfWarpBaseAddress + N

         HalfWarpBaseAddress应该排列成16 * sizeof ( type ) 字节，例如16 * sizeof ( type ) 的倍数。

 任何一个驻留在全局内存BaseAddress的变量的地址，或者一个来自 D.5 或 E.8 部分内存分配规则返回的地址，问题被排列成至少256个字节，以此来满足内存队列的约束，

建议：对于整个warp满足这个需求，而不是分开的，半个半个的。

            因为未来的设备将默认为必要的要求。

示例：联合的 64-bit 访问会比联合的 32-bit 访问内存带宽低一点，

            联合的128-bit访问会比联合的32-bit访问内存带宽要低很多。

  

 

一个公共的全局内存访问样式是，当每个带有线程ID tid 的线程访问位于一个数组的一个元素时，元素的地址位于类型 type* 的 BaseAddress，使用以下地址

如果 type 的结构 > 16字节，它应该被分成几个满足要求的结构，并且数据应该在内存中被划分成关于这些结构的几个数组，而不是一个类型 type* 的单一数组。

另一个公共的全局内存访问样式是，

当带有索引 (tx, ty) 的每条线程访问地址位于类型 type* 的BaseAddress 和宽度 width 的2D数组的一个元素使用以下地址：

a. 块线程的宽度是一半warp大小的倍数

b. width 是16 的倍数

特别是，这意味着，宽度不是16的倍数的数组将被更高效地访问，如果它实际上分配是宽度被传入到最接近16的倍数而且，它的列因此被填充了。

cudaMallocPitch() 和 cuMemAllocPitch() 函数和在D.5和E.8 部分描述的相关内存拷贝的函数，使开发人员能够编写非硬件独立的代码，来分配遵循这些约束的数组。