CUDA_全局内存及访问优化

全局内存

GPU全局内存，CPU和GPU均可以进行读写操做。任何设备均可以经过PCI-E总线对其进行访问，GPU之间不经过CPU，直接将数据从一块GPU卡上的数据传输到另外一块GPU上。

点对点的特性实在DUDA4.x SDK中引入。只对特定平台进行支持（特斯拉硬件经过TCC驱动模型可以支持windows7和windows Vista平台，对于linux或windowsXP平台，消费机GPU卡和特斯拉卡都支持）。

CPU主机端处理器能够经过如下三种方式对GPU上的内存进行访问：

• 显式地阻塞传输；
• 显式地非阻塞传输；
• 隐式地使用零拷贝内存复制。

一旦数据进入到GPU，主要问题就成了如何在GPU中进行高效访问。经过建立一个每十次计算只需一次访存的模式，内存延迟能明显的被隐藏，但前提是对全局内存的访问必须是以合并的方式进行访问。

对全局内存的访问是否知足合并访问条件是对CUDA程序性能影响最明显的因素之一。

合并访问--全局存储器访问优化

全部线程访问连续的对齐的内存块。

若是咱们对内存进行一对一连续对齐访问，则每一个线程的访问地址能够合并起来，只需一次存储食物便可解决问题。假设咱们访问一个单精度或者整型值，每一个线程将访问一个4字节的内存块。内存会基于线程束的方式进行合并（老式的G80硬件上使用半个线程束），也就是说访问一次内存将获得32*4=128个字节的数据。

合并大小支持32字节、64字节、128字节，分贝标识线程束中每一个线程一个字节、16位以及32位为单位读取数据，但前提是访问必须连续，而且以32字节位基准对其。

将标准的cudaMalloc替换为cudaMallocPitch，能够分配到对齐的内存块。

extern __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height);

该方法的第一个参数表示指向设备内存指针的指针，第二个参数表示指向对齐以后每行真实字节数的指针，第三个参数为须要开辟的数据的宽度，单位为字节，最后一个参数为数组的高度。

合并访问条件要求同一warp或者同一half-warp中的线程要按照必定字长访问通过对齐的段。

不一样设备中合并访问的具体要求：

• 计算能力1.0、1.1设备上，一个half-warp中的第k个线程必须访问段里面的第k个字，而且half-warp访问的段的地址必须对齐到每一个线程访问的字长的16倍。只支持对字长32bit、64bit、128bit的数据的合并访问。
• 在1.2及更高能力的设备上，合并访问要求大大放宽，支持字长为8bit（对应段长32Byte）、16bit（对应段长64Byte）、32bit/64bit/128bit（对应段长128Byte）的数据进行合并访问。

下面描述1.2/1.3能力硬件的一个half-warp是如何完成一次合并访问的。

• 首先，找到有最低线程号活动线程（前half-warp中的线程0，或者后half-warp中的线程16）请求访问的地址所在段。对于8bit数据来讲，段长为32Byte，对于16bit数据来讲段长为64Byte，对于3二、6四、128bit数据来讲段长为128Byte。
• 而后，找到所请求访问的地址也在这个段内的活动线程。若是全部线程访问的数据都处于段的前半部分或者后半部分，那么还能够减小一次传输的数据大小。例如，若是一个段的大小为128Byte，但只有上半部分或下半部分被使用了，那么实际传输的数据大小就能够进一步减少到64Byte，同理，对于64Byte的段的合并传输，在只有前半或者后半被使用的状况下也能够继续减少到32Byte。
• 进行传输，此时，执行访存指令的线程将处于不活动状态，执行资源被释放供SM中处于就绪态的其余warp使用。
• 重复上述过程，知道half-warp全部线程均访问结束。

须要注意的是，经过运行时API(如cudaMalloc())分配的存储器，已经能保证其首地址至少会按256Byte进行对齐。所以，选择合适的线程块大小（例如16的整数倍），能使half-warp的访问请求按段长对齐。使用__align__(8)和__align__(16)限定符来定义结构体，可使对结构体构成的数组进行访问时可以对齐到段。

访问时段不对齐或者间隔访问都会要成有效带宽的大幅度下降。对于间隔访问显存的状况，能够借助shared memory来实现。

全局内存分配

当使用CUDA运行时时，设备指针与主机指针类型均为void*。

动态内存分配

大多数CUDA中的全局内存经过动态分配获得，使用cuda运行时，经过如下函数分别进行全局内存的分配和释放。

cudaError_t cudaMalloc(void **, size_t);
cudaError_t cudaFree(void);

对应的驱动程序API函数为：

CUresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);
CUresult CUDAAPI cuMemFree(CUdeviceptr dptr);

分配全局内存成本较大，CUDA驱动程序实现了一个CUDA小型内存请求的子分配器（suballocator），可是若是这个suballocator必须建立一个新的内存块，这须要调用操做系统的一个成本很高的内核模式驱动程序。若是这种状况发生，CUDA驱动程序必须与GPU同步，这可能会中断CPU、GPU的并发，所以，在性能要求很高的代码中避免分配或释放全局内存时一个较好的作法。

静态内存分配

经过使用__device__关键字标记在内存声明中进行标记便可。这一内存是由cuda驱动程序在模块加载时分配的。

运行时API：

cudaError_t cudaMemcpyToSymbol(
	char *symbol,
	const void *src,
	size_t count,
	size_t offset=0,
	enum cudaMemcpyKind kind=cudaMemcpyHostToDevice
);

cudaError_t cudaMemcpyFromSymbol(
	void *dst,
	char *symbol,
	size_t count,
	size_t offset,
	enum cudaMemcpyKind kind=cudaMemcpyDeviceToHost
);

cuda运行时应用程序能够经过调用函数cudaGetSymbolAddress()查询关联到静态分配的内存上的指针。

cudaError_t cudaGetSymbolAddress(void **devPtr, char *symbol);

驱动程序API：

CUresult CUDAAPI cuModuleGetGlobal(CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name);

该函数返回基指针和对象大小。若是咱们不须要大小，能够在bytes参数传入NULL。

指针查询

cuda跟踪全部内存分配，并提供API使应用程序能够查询CUDA中的全部指针。函数库和插件能够在基础之上使用不一样的处理策略。

struct cudaPointerAttributes{
  	enum cudaMemoryType memoryType;
  	int device;
  	void *devicePointer;
  	void *hostPointer;
}