CUDA C编程5 - CUDA内存模型

系列文章目录

---

---

前言

最近在温习CUDA C 全局内存的知识，这里对关键知识点进行总结并分享给大家。

较差的全局内存访问方式是造成内存负载效率大幅下降的原因之一。

这里主要介绍以下几点内容：

1. 剖析核函数与全局内存的联系及其对性能的影响；
2. 介绍全局内存访问模式
3. 介绍如何通过核函数高效的利用全局内存

---

一、CUDA内存模型

现代加速器中，内存管理对高性能计算有着极大的影响。

大多数工作负载收到加载和存储速度限制，因此低延迟、高带宽内存对性能十分有利。

CUDA内存模型结合了主机和设备的内存系统，展现了完整的内存层次结构，这样我们可以显式的控制数据布局以达到优化性能的目的。

现代的计算机不断改进低延迟低容量的内存层次结构来优化性能，典型的内存层次结构如下图所示：
CPU和GPU的主存采用DRAM(动态随机存取存储器)，低延迟内存（缓存、寄存器）使用SRAM（静态随机存取存储器）。内存层次结构中最大且最慢的级别通常使用磁盘或闪存驱动来实现。

如果数据被处理器频繁使用，该数据应该被保存在低延迟、低容量的存储器中；如果数据被存储以备后用，数据就该被存储在高延迟、大容量的存储器中。

GPU和CPU内存模型的主要区别是，CUDA编程模型能将内存层次结构更好的让我们能显示的控制它的行为。

对于我们这种程序猿来说，存储器有两种类型：可编程与不可编程的。

在CPU内存层次结构中，L1和L2缓存是不可编程的；

CUDA内存模型中可编程内存有：

1. 寄存器
2. 共享内存
3. 本地内存
4. 常量内存
5. 纹理内存
6. 全局内存

上图中的内存空间的层次结构中，每种都有不同的作用域、生命周期和缓存行为。
核函数中的线程有自己私有的本地内存；一个线程块有自己的共享内存，同一线程块中所有线程都可见，内容持续线程块的整个生命周期；所有线程都可以访问全局内存；所有线程都能访问的只读内存空间：常量内存空间和纹理内存空间。

这里需要特别提示的是：纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。

对于一个应用程序，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

1. 寄存器

寄存器是GPU上运行速度最快的内存空间。
核函数中声明的没有其他修饰符的自变量（个人理解，就是普通的函数内局部变量）通常存储在寄存器中。如果用于引用该数组的索引是常量且能在编译时确定，该数组也存储在寄存器中。

寄存器变量对于每个线程时私有的，核函数通常使用寄存器来保存需要频繁访问的线程私有变量。寄存器变量与核函数的生命周期相同，核函数执行结束，那么就不能对寄存器变量进行访问。

在不同的GPU中，每个线程拥有的寄存器是有限的，本人的超薄本MX250寄存器数量可用如下命令查看：

cudaDeviceProp stDeviceProp;
cudaGetDeviceProperties(&stDeviceProp, nDeviceId);
printf("Total number of registers available per block: %d\n",stDeviceProp.regsPerBlock);

在核函数使用较少的寄存器，那么在SM上有更多的常驻线程块。每个SM上并发线程块越多，使用率和性能就越高。

可以使用nvcc编译器选项检查核函数使用的硬件资源，如：寄存器数量，在windows 2017上属性页面设置如下：
![在这里插入图片描述](https://img-blog.csdnimg.cn/f69f64e8a4574bc2a72762f94994f5f0.png?x-oss-process=image/watermark,type_d3F5LXplbmhlaQ,shadow_50,text_Q1NETiBARFVfWVVMSU4=,size_20,color_FFFFFF,t_70,g_se,x_16

VS2017 寄存器数量编译输出：
一旦核函数使用超过硬件限制数量的寄存器，则会使用本地内存替代多占用的寄存器。这种寄存器溢出给性能带来不利影响。nvcc编译器使用启发式策略来最小化寄存器使用，以避免寄存器溢出。

可以使用maxrregcount编译器选项来控制核函数使用的寄存器最大数量，vs2017配置选项如上图项目属性页面。

2. 本地内存

如上所述，本该进入寄存器但是因空间不足无法进入的变量将溢出到本地内存中。对于计算力在2.0及以上的GPU,本地内存数据存储在每个SM的一级缓存和每个设备的二级缓存中。

3. 共享内存

共享内存存放由修饰符 $\_\_shared\_\_$修饰的变量，与本地内存和全局变量相比，它具有更高的带宽和耕地的延迟。类似CPU的一级缓存，它是可编程的。

每个SM有一定数量的由线程块分配的共享内存。同时要注意，必须非常小心并不要过度使用共享内存，否则将不经意间限制活跃线程束的数量。

共享内存在核函数内声明，生命周期与线程块相同。

共享内存是线程块内线程间相互通信的基本方式。访问共享内存需要使用同步方法：

void __syncthreads();

在核函数中调用上面函数，可以避免潜在的数据冲突，但是也会通过频繁强制SM到空闲状态影响性能。

SM中的一级缓存和共享内存可通过如下函数进行动态配置

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);

动态配置类型：

enum __device_builtin__ cudaFuncCache
{
    cudaFuncCachePreferNone   = 0,    /**< Default function cache configuration, no preference */
    cudaFuncCachePreferShared = 1,    /**< Prefer larger shared memory and smaller L1 cache  */
    cudaFuncCachePreferL1     = 2,    /**< Prefer larger L1 cache and smaller shared memory */
    cudaFuncCachePreferEqual  = 3     /**< Prefer equal size L1 cache and shared memory */
};

4. 常量内存

常量内存使用修饰符 $\_\_constant\_\_$, 常量变量必须在全局空间内核所有核函数之外进行声明。

常量内存是静态声明的，对同一编译单元中所有核函数可见。

核函数只能从常量内存中读取数据，常量内存必须在主机端使用下面函数来初始化：

extern __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice));

大多数情况下，cudaMemcpyToSymbol函数时同步的。

线程束中所有线程从相同的内存地址中读取数据时，常量内存表现最好。如果线程束里每个线程都从不同的地址空间读取数据，并且只读依次，那么常量内存中就不是最佳选择。

5. 纹理内存

纹理内存是一种通过指定的只读缓存访问的全局内存。纹理内存是对二维空间局部性的优化，因此纹理内存在线程束里使用（访问二维数据）可以使线程性能达到最优。注意只读缓存拥有滤波硬件的支持，可将浮点插入作为读过程的一部分来执行。但是不是所有的情况都适合使用纹理内存，在有些应用程序中，使用纹理内存可能比全局内存更慢。（照搬参考资料中的说法，以后遇到案例再另行说明吧）

6. 全局内存

全局内存是GPU中最大、延迟最高但是最长使用的内存。
一个全局内存变量可以是静态声明或动态声明：
（1）静态声明需要使用修饰符：$\_\_device\_\_$;
（2）动态声明，即在CPU中使用$cudaMalloc$函数分配的全局内存，使用完后需要调用$cudaFree$函数释放全局内存。

全局内存分配空间存在于应用程序的整个生命周期，并且可以被所有核函数中所有线程访问。这也是多线程访问全局内存时需要特别注意的点，一旦使用不当，很容易产生脏数据，即不同线程块内线程可能同时修改全局内存中同一地址的数据，很明显将产生未定义的程序行为。

全局内存常驻于设备内存中，可通过32字节、64字节或128字节的内存事务进行访问。这些内存事务必须自然对齐，即首地址必须是32字节、64字节或128字节的倍数。优化内存事务对于获得最优性能来说是格外重要的。（内存事务包括软件内存事务和硬件内存事务，软件内存事务主要使用原子对象和冲突判决器，硬件内存事务这里就不介绍了，个人也不熟悉，感兴趣的同学可以自行搜索）。

当线程束执行内存加载或存储时，需要满足的传输数量通常取决于以下两个因素：
（1）跨线程的内存地址分布；
（2）每个事务内存地址的对齐方式；

一般情况下，用来满足内存请求的事务越多，未使用的字节被传输回的可能性就越高，这就造成了数据吞吐率的降低。

对于一个给定的线程束内存请求，事务数量和数据吞吐率是由设备的计算能力确定的。缓存的内存事务利用数据局部性来提高数据吞吐率。

什么是数据局部性呢？一般来说，我们编写的程序不会在某一时间点访问任意的数据或运行任意的代码。应用程序往往遵循局部性原则，即在任意时间点访问相对较小的局部地址空间。有两种不同类型的局部性：
（1）时间局部性
即如果一个数据位置被使用，那么该数据在较短的时间内很可能会被再次使用，随着时间流逝，该数据被使用的可能性也逐渐降低。
（2）空间局部性
如果一个内存位置被访问，则附近的位置也可能会被访问。

7. GPU缓存

与CPU缓存相同，GPU缓存也是不可编程的内存。
在GPU中有四种缓存：
（1）一级缓存
（2）二级缓存
（3）只读常量内存
（4）只读纹理内存

每个SM（SM, Stream Multiprocessor - 流式多处理器）都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。

在CPU上，内存的加载和存储都可被缓存，但在GPU上，只有内存加载操作可以被缓存，内存存储操作不能；

每个SM也有一个只读常量缓存和只读纹理缓存，用于在设备内存中提高读取性能。

上面说了那么多，下面应用参考资料中的图表来总结CUDA变量声明和相应存储器、作用域、生命周期和修饰符。

下表总结了存储器的重要特征

二、静态全局内存示例

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <device_functions.h>


//declare a static global variable
__device__ float fDevData;

//define a device(kernel) function to modify the value of the static global variable
__global__ void Device_CheckGlobalVariable() {
	//print origin value
	printf("Device:     the value of the global variable is %f\n", fDevData);

	//modify the static global variable
	fDevData += 2.0f;
}

//main function - CPU
int main() {
	float value = 3.14f;

	//copy CPU data to the device data
	//note: mustn't add '&' for the static global variable 'fDevData', otherwise the copy operate will loss it's function
	cudaMemcpyToSymbol(fDevData, &value, sizeof(float));
	cudaDeviceSynchronize();
	printf("Host:    copyied %f to the global variable\n", value);

	//call the device function
	Device_CheckGlobalVariable << <1, 1 >> > ();

	cudaDeviceSynchronize();

	// copy the global variable back to the host
	// note: mustn't add '&' for the static global variable 'fDevData', otherwise the copy operate will loss it's function
	cudaMemcpyFromSymbol(&value, fDevData, sizeof(float));

	//cudaDeviceSynchronize();

	printf("Host:    the value changed by the kernel to %f\n", value);

	cudaDeviceReset();

	system("pause");
	return 0;
}

上述静态全局变量通过$cudaMemcpyToSymbol()$函数将CPU中变量值赋值给设备全局变量，又通过$cudaMemcpyFromSymbol()$函数将设备全局变量值拷贝到CPU变量中，这里有一点要特别注意，不要给静态全局变量添加取地址符号“&”，这样会造成拷贝失效，即无法给设备全局变量赋值或者拷贝设备全局变量的值到CPU变量中。这是因为设备全局变量只是在GPU设备上表示物理位置的符号，不是简单的通过取地址符号获得设备全局变量的地址。

但是可通过$cudaGetSymbolAddress()$函数获得与设备符号相关的设备全局变量的物理地址，然后就可以使用$cudaMemcpy（）$函数来给设备变量赋值，或从设备变量中取值。

注意，虽然设备函数(GPU)与主机函数（CPU）在相同的文件中，但是他们的变量不能直接访问，要通过$cudaMemcpy$函数来传递值得方式。但是有一个例外，可以直接从主机引用GPU内存，那就是CUDA固定内存。主机函数和设备函数都可以通过简单的指针引用直接访问固定内存。（固定内存的内容后面会更新，敬请期待！）

正确运行结果：

cudaMemcpyToSymbol、cudaMemcpyFromSymbol函数中对设备全局变量使用取地址符号，运行的错误结果如下：
从错误结果可知，上面的两个函数对设备全局变量取地址后，函数功能失效了。

---

总结

这次的笔记还是挺重要的，对以后cuda编程思想是基础中的重点，只有很好地理解了内存模型，后面才能写出更加高性能的cuda程序，也能更快的定位问题，基础很重要！！！

---

参考资料

《CUDA C编程权威指南》

---

附录

1. 设备属性信息查看C/C++源码

#include <stdio.h>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>
#include <device_functions.h>

int main()
{
	int nDeviceId = 0;
	cudaDeviceProp stDeviceProp;
	cudaGetDeviceProperties(&stDeviceProp, nDeviceId);

	printf("Device %d: %s\n", nDeviceId, stDeviceProp.name);
	printf("Number of multiprocessors(SM): %d\n", stDeviceProp.multiProcessorCount);
	printf("Total amount of constant memory: %4.2f KB\n",
		stDeviceProp.totalConstMem / 1024.0);
	printf("Total amount of shared memory per block: %4.2f KB\n",
		stDeviceProp.sharedMemPerBlock / 1024.0);
	printf("Total number of registers available per block: %d\n",
		stDeviceProp.regsPerBlock);
	printf("Warp size: %d\n", stDeviceProp.warpSize);
	printf("Max number of threads per block: %d\n",
		stDeviceProp.maxThreadsPerBlock);
	printf("Max number of threads per multiprocessor: %d\n",
		stDeviceProp.maxThreadsPerMultiProcessor);
	printf("Max number of warps per multiprocessor: %d\n",
		stDeviceProp.maxThreadsPerMultiProcessor / stDeviceProp.warpSize);
	
	system("pause");
	return 0;
}