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，对于32、64、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;
}

CUDA_全局内存及访问优化

原文：https://www.cnblogs.com/TonvyLeeBlogs/p/13951351.html