CUDA编程（六）CUDA的内存组织

1.内存组织简介：

现代计算机中的内存往往存在一种组织结构（hierarchy）。在这种结构中，含有多种类 型的内存，每种内存分别具有不同的容量和延迟（latency，可以理解为处理器等待内存数 据的时间）。一般来说，延迟低（速度高）的内存容量小，延迟高（速度低）的内存容量大。 当前被处理的数据一般存放于低延迟、低容量的内存中；当前没有被处理但之后将要被处理的大量数据一般存放于高延迟、高容量的内存中。相对于不用分级的内存，用这种分级的内存可以降低延迟，提高计算效率。

总结，内存两个重要参数容量和延迟，一般，容量大，延迟高。

global，constant，texture，register，local，shares   memory。

2.CUDA中不同类型的内存

全局内存：

指核函数中所有线程都能够访问其中的数据。全局内存（Global Memory）是指在CUDA编程中用于在GPU设备上存储数据的主要内存区域之一。全局内存是GPU上全局可访问的内存，它可以被所有线程（也称为CUDA核函数）访问。

特点，全局性（所有线程都能访问），持久性，访问慢（相比寄存器和缓存），可与主机内存数据传输（一般用cudaMemcpy搬运数据，注意数据传输方向），需要程序员显示分配和释放（cudaMalloc，cudaFree）。

全局内存由于没有存放在GPU的芯片上，因 此具有较高的延迟和较低的访问速度。然而，它的容量是所有设备内存中最大的。其容量基本上就是显存容量。

全局内存的主要角色是为核函数提供数据，并在主机与设备及设备与设备之间传递数据。

首先，我们用cudaMalloc函数为全局内存变量分配设备内存。然后，可以直接在核函 数中访问分配的内存，改变其中的数据值。

全局内存可读可写。一个网格的所有线程都可以访问（读 或写）传入核函数的设备指针所指向的全局内存中的全部数据。

全局内存的生命周期（lifetime）不是由核函数决定的，而是由主机端决定的。从cudaMalloc开始，到cudaFree结束。

在处理逻辑上的两维或三维问题时，可以用cudaMallocPitch和cudaMalloc3D函数分配内存，用cudaMemcpy2D 和 cudaMemcpy3D 复制数据，释放时依然用cudaFree 函数。

以上所有的全局内存都称为线性内存（linearmemory）。在CUDA中还有一种内部构 造对用户不透明的（nottransparent）全局内存，称为CUDAArray。CUDAArray使用英伟 达公司不对用户公开的数据排列方式，专为纹理拾取服务。

全局内存变量都是动态地分配内存的。在CUDA中允许使用静态全局 内存变量，其所占内存数量是在编译期间就确定的。而且，这样的静态全局内存变量必须 在所有主机与设备函数外部定义，所以是一种“全局的静态全局内存变量”。这里，第一 个“全局”的含义与C++中全局变量的含义相同，指的是对应的变量对从其定义之处开始、 一个翻译单元内的所有设备函数直接可见。如果采用所谓的分离编译（separatecompiling）， 还可以将可见范围进一步扩大。

静态全局内存变量由以下方式在任何函数外部定义：

__device__ T x; // 单个变量

__device__ T y[N]; // 固定长度的数组

其中，修饰符__device__说明该变量是设备中的变量，而不是主机中的变量；T是变量的 类型；N是一个整型常数。

在核函数中，可直接对静态全局内存变量进行访问，并不需要将它们以参数的形式传给核函数。不可在主机函数中直接访问静态全局内存变量，但可以用cudaMemcpyToSymbol函数和 cudaMemcpyFromSymbol 函数在静态全局内存与主机内存之间传输数据。

cudaError_t cudaMemcpyToSymbol{
    const void* symbol,  //静态全局内存变量名
    const void* src,     //主机内存缓冲区指针
    size_t count,        //复制的字节数
    size_t offset = 0,   //从symbol对应设备地址开始偏移的字节数
    cundaMemcpyKind kind = cudaMemcpyHostToDevice //可选参数
}

cudaError_t cudaMemcpyFromSymbol{
     void* dst,  //主机内存缓冲区指针
    const void* symbol,     //静态全局内存变量名
    size_t count,        //复制的字节数
    size_t offset = 0,   //从symbol对应设备地址开始偏移的字节数
    cundaMemcpyKind kind = cudaMemcpyDeviceToHost //可选参数
}

这两个函数的参数symbol可以是静态全局内存变量的变量名，也可以是常量内存变量的变量名。

看个例子：

#include "error.cuh"
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include<stdlib.h>


//全局变量存放在GPU全局内存中的变量，核函数中可以直接调用，不需要向核函数传参。
__device__ int d_x = 1;
__device__ int d_y[2];

void __global__ my_kernel(void) {
	d_y[0] += d_x;
	d_y[1] += d_x;
	printf("d_x = %d,d_y[0] = %d, d_y[1] = %d.\n", d_x, d_y[0], d_y[1]);
}

int main(void) {
	int h_y[2] = { 10,20 };
	CHECK(cudaMemcpyToSymbol(d_y, h_y, sizeof(int) * 2));

	my_kernel << <1, 1 >> > ();

	CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));
	//sizeof(int) * 2，数据大小，这么写安全。

	printf("h_y[0] = %d,h_y[1] = %d，sizeof(int) * 2 = %d.\n", h_y[0], h_y[1], sizeof(int) * 2);


	return 0;
}

有注释，cudaMemcpyToSymbol将主机数据h_y传到设备上的全局内存变量d_y上。

cudaMemcpyFromSymbol将设备上的全部内存变量传到主机上的变量h_y来。

程序输出：

而C++中的全局变量，存放在全局存储区域，程序启动，到结束，与程序声明周期一样长。

常量内存：

常量内存（constantmemory）是有常量缓存的全局内存，数量有限，一共仅有64KB。 它的可见范围和生命周期与全局内存一样。

不同的是，常量内存仅可读、不可写。由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程（一个线程块中相邻的32个线程）要读取相同的常量内存数据。

一个使用常量内存的方法是在核函数外面用__constant__ 定义变量，并用前面介绍 的CUDA运行时API函数cudaMemcpyToSymbol将数据从主机端复制到设备的常量内存后供核函数使用。当计算能力不小于2.0时，给核函数传递的参数（传值，不是像全局变量 那样传递指针）就存放在常量内存中，但给核函数传递参数最多只能在一个核函数中使用 4 KB常量内存。

在数组相加的例子中，核函数的参数 const int N 就是在主机端定义的变量，并通过传值的方式传送给核函数中的线程使用。在核函数中的代码段 if (n < N)中，这个参数N就被每一个线程使用了。所以， 核函数中的每一个线程都知道该变量的值，而且对它的访问比对全局内存的访问要快。除 给核函数传递单个的变量外，还可以传递结构体，同样也是使用常量内存。结构体中可以定义单个的变量，也可以定义固定长度的数组。

纹理内存和表面内存

纹理内存（texture memory）和表面内存（surface memory）类似于常量内存，也是一 种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读（表面内存也可 写）。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。

对于计算能力不小于3.5的GPU来说，将某些只读全局内存数据用__ldg()函数通过 只读数据缓存（read-onlydatacache）读取，既可达到使用纹理内存的加速效果，又可使代码简洁。

该函数的原型为

T __ldg(const T* address);

其中，T是需要读取的数据的类型；address是数据的地址。对帕斯卡架构和更高的架构来说，全局内存的读取在默认情况下就利用了__ldg()函数，所以不需要明显地使用它。

寄存器

在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器（register）中。 核函数中定义的不加任何限定符的数组有可能存放于寄存器中，但也有可能存放于 局部内存中。另外，以前提到过的各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中。在核函数中访问这些内建变量是很高效的。

在数组求和的例子中，我们在核函数中有如下语句：

const int n = blockDim.x * blockIdx.x + threadIdx.x;

这里的n就是一个寄存器变量。寄存器可读可写。上述语句的作用就是定义一个寄存器变 量n并将赋值号右边计算出来的值赋给它（写入）。

z[n] = x[n] + y[n];

寄存器变量n的值被使用（读出）。

寄存器变量仅仅被一个线程可见。也就是说，每一个线程都有一个变量n的副本。虽然在核函数的代码中用了这同一个变量名，但是不同的线程中该寄存器变量的值是可以不同的。每个线程都只能对它的副本进行读写。寄存器的生命周期也与所属线程的生命周期 一致，从定义它开始，到线程消失时结束。

寄存器内存在芯片上（on-chip），是所有内存中访问速度最高的，但是其数量也很有限。

一个寄存器占有32bit（4字节）的内存。所以，一个双精度浮点数将使用两个 寄存器。这是在估算寄存器使用量时要注意的。

局部内存

目前还没有用过局部内存（localmemory），但从用法上看，局部内存和寄存器几乎一 样。

核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。

寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。 这种判断是由编译器自动做的。

对于数组相加例子中的变量n来说，作者可以肯定它在寄存器中，而不是局部内存中，因为核函数所用寄存器数量还远远没有达到上限。

虽然局部内存在用法上类似于寄存器，但从硬件来看，局部内存只是全局内存的一部分。所以，局部内存的延迟也很高。每个线程最多能使用高达512KB的局部内存，但使用过多会降低程序的性能。

共享内存

没有使用过共享内存（sharedmemory）。共享内存和寄存器类似，存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。

不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。也就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本。共享内存的主要作用是减少对全局内存的访问，或者改善对全局内存的访问模式。

L1和L2缓存

从费米架构开始，有了SM层次的L1缓存（一级缓存）和设备（一个设备有多个SM） 层次的L2缓存（二级缓存）。它们主要用来缓存全局内存和局部内存的访问，减少延迟。

从硬件的角度来看，开普勒架构中的L1缓存和共享内存使用同一块物理芯片；在麦克 斯韦架构和帕斯卡架构中，L1缓存和纹理缓存统一起来，而共享内存是独立的；在伏特架 构和图灵架构中，L1缓存、纹理缓存及共享内存三者统一起来。

从编程的角度来看，共享内存是可编程的缓存（共享内存的使用完全由用户操控），而L1和L2缓存是不可编程的 缓存（用户最多能引导编译器做一些选择）。

对某些架构来说，还可以针对单个核函数或者整个程序改变L1缓存和共享内存的比 例。

3.SM及其占有率

流式多处理器SM

一个 GPU是由多个SM构成的。

GPU（图形处理单元）架构中，"SM" 是指 "Streaming Multiprocessor"，中文译为"流处理多处理器"。SM 是现代GPU的核心组成部分之一，用于执行并行计算任务。每个GPU设备通常包含多个 SM，每个 SM 可以并行执行多个线程，以加速计算任务。

开发者通常将计算任务分配给线程块，然后由 SM 来执行这些线程块的任务。

一个SM包含如下资源：

寄存器，共享内存，常量内存的缓存，纹理和表面内存的缓存，L1缓存，

两个（计算能力6.0）或4个（其他计算能力）线程束调度器（warpscheduler），用于在不同线程的上下文之间迅速地切换，以及为准备就绪的线程束发出执行指令。

执行核心，包括：

– 若干整型数运算的核心（INT32）

– 若干单精度浮点数运算的核心（FP32）

– 若干双精度浮点数运算的核心（FP64）

– 若干单精度浮点数超越函数（transcendentalfunctions）的特殊函数单元（Special Function Units，SFUs）

– 若干混合精度的张量核心（tensorcores，由伏特架构引入，适用于机器学习中的 低精度矩阵计算）。

其中，"超越函数"（Transcendental Function）是数学中的一类特殊函数，通常用于描述无理数和复杂数的数学运算。如指数，对数，三角，双曲，伽马，贝塞尔等函数。

SM的占有率

因为一个SM中的各种计算资源是有限的，那么有些情况下一个SM中驻留的线程数 目就有可能达不到理想的最大值。

此时，我们说该SM的占有率小于100%。获得100%的 占有率并不是获得高性能的必要或充分条件，但一般来说，要尽量让SM的占有率不小于 某个值，比如25%，才有可能获得较高的性能。

当并行规模较小时，有些SM可能就没有被利用，占有率为零。这是导致程序性能低下的原因之一。当并行规模足够大时，也有可能得到非100%的 占有率，这就是下面要讨论的情形。

要分析SM的理论占有率（theoretical occupancy），还需要知道两个指标：

• 一个SM中最多能拥有的线程块个数为Nb = 16（开普勒架构和图灵架构）或者 Nb = 32（麦克斯韦架构、帕斯卡架构和伏特架构）；

• 一个SM中最多能拥有的线程个数为Nt = 2048（从开普勒架构到伏特架构)或者Nt = 1024（图灵架构）。

寄存器和共享内存使用量很少的情况。此时，SM的占有率完全由执行配置中的线程 块大小决定。关于线程块大小，读者也许注意到我们之前总是用128。这是因为，SM中 线程的执行是以线程束为单位的，所以最好将线程块大小取为线程束大小（32个线 程）的整数倍。例如，假设将线程块大小定义为100，那么一个线程块中将有3个完 整的线程束（一共96个线程）和一个不完整的线程束（只有4个线程）。在执行核函数中的指令时，不完整的线程束花的时间和完整的线程束花的时间一样，这就无形中浪费了计算资源。所以，建议将线程块大小取为32的整数倍。在该前提下，任何 不小于Nt/Nb而且能整除Nt的线程块大小都能得到100%的占有率。根据我们列出的数据，线程块大小不小于128时开普勒架构能获得100%的占有率；线程块大小不小于64时其他架构能获得100%的占有率。

有限的寄存器个数对占有率的约束情况。对于表6.2中列出的所有计算能 力，一个SM最多能使用的寄存器个数为64K(64X1024)。除图灵架构外，如果我们希望在一个SM中驻留最多的线程（2048个,2X1024），核函数中的每个线程最多只能用32个寄存器。当每个线程所用寄存器个数大于64时，SM的占有率将小于50%；当每个线程所用寄存器个数大于128时，SM的占有率将小于25%。对于图灵架构，同样的占有率允许使用更多的寄存器。

有限的共享内存对占有率的约束情况。因为共享内存的数量随着计算能力的上升没有显著的变化规律，所以我们这里仅针对一个3.5的计算能力进行分析，对其他计算能力可以类似地分析。如果线程块大小为128(2^7)，那么每个SM要激活16(2^4)个线程块才能有2048(2^11)个线程，达到100%的占有率。此时，一个线程块最多能使用3KB的共享内存。在不改变线程块大小的情况下，要达到50%的占有率，一个线程块最多能使用6KB的共享内存；要达到25%的占有率，一个线程块最多能使用12KB共享内存。最后，如果一个线程块使用了超过48KB的共享内存，会直接导致核函数无法运行。对 其他线程块大小可类似地分析。

•以上单独分析了线程块大小、寄存器数量及共享内存数量对SM占有率的影响。一般情况下，需要综合以上三点分析。在CUDA工具箱中，有一个名 为CUDA_Occupancy_Calculator.xls的Excel文档，可用来计算各种情况下 的SM占有率，感兴趣的读者可以去尝试使用。 值得一提的是，用编译器选项--ptxas-options=-v可以报道每个核函数的寄存器使用数量。CUDA还提供了核函数的__launch_bounds__()修饰符和--maxrregcount=编译选项来让用户分别对一个核函数和所有核函数中寄存器的使用数量进行控制。

4.用CUDA运行时的API函数查询设备

话不多数，直接看代码：

#include "error.cuh"
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include<stdlib.h>

int main(int argc, char* argv[]) {
	int device_id = 0;
	if (argc > 1) device_id = atoi(argv[1]);
	CHECK(cudaSetDevice(device_id));

	cudaDeviceProp prop;
	CHECK(cudaGetDeviceProperties(&prop, device_id));
	
	printf("Deviceid: %d\n", device_id);
	
	printf("Device name: %s\n", prop.name);

	printf("Compute capability: %d.%d\n",prop.major, prop.minor);
	
	printf("Amountof globalmemory: %g GB\n",prop.totalGlobalMem / (1024.0 * 1024 * 1024));
	
	printf("Amountof constantmemory: %g KB\n", prop.totalConstMem / 1024.0);
	
	printf("Maximum gridsize: %d %d %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    
	printf("Maximum block size: %d %d %d\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
	
	printf("Numberof SMs: %d\n",prop.multiProcessorCount);
   
	printf("Maximum amountof sharedmemory per block: %g KB\n", prop.sharedMemPerBlock / 1024.0);
	
	printf("Maximum amountof sharedmemory per SM:%g KB\n",prop.sharedMemPerMultiprocessor / 1024.0);
	
	printf("Maximum numberof registers per block:%d K\n", prop.regsPerBlock / 1024);
	
	printf("Maximum numberof registers per SM: %dK\n", prop.regsPerMultiprocessor / 1024);
    
	printf("Maximum numberof threads perblock: %d\n", prop.maxThreadsPerBlock);
    
	printf("Maximum numberof threads perSM: %d\n", prop.maxThreadsPerMultiProcessor);
	
	return 0;

}

然后就是结果

这些参数可以对比着看，后面在使用时很重要。

此外，还有一个知识点（来自GPT）

在 C++ 的 main 函数中，int argc 和 char* argv[] 是标准的参数，它们通常用于接收命令行参数传递给程序的信息。这两个参数的含义如下：

1. int argc：这是一个整数参数，表示命令行参数的数量，即程序运行时传递的参数的总数。这个参数至少为 1，因为第一个参数通常是程序的名称（可执行文件的路径），后面可以包括其他参数。

2. char* argv[]：这是一个字符指针数组，通常称为参数数组（Argument Array）。它包含了指向每个命令行参数的字符串指针。argv[0] 指向程序的名称，argv[1]、argv[2]、等等依次指向其他传递的参数。每个参数都是一个以空格分隔的字符串。

通过这两个参数，你可以在程序中访问并解析命令行参数，以根据需要进行不同的操作。这对于控制程序的行为、传递配置选项或文件路径以及与用户进行交互等方面非常有用。

一般在linux上常用。