CUDA C++ / 内存管理以及优化

内存

1 内存读写速度

线程寄存器读写：1个时钟周期延迟

线程本地内存读写：非常慢

块的共享内存读写：1个时钟周期延迟，但是可能冲突

网格全局内存读写：500个时钟周期延迟，联合访问时会有隐含延迟。

网格的常量内存和纹理内存读取：500个时钟周期延迟。但是有缓存

2 内存模型

2.1 寄存器

修饰符：核函数中声明，没有修饰符。

共享范围：线程私有

速度：块。一个时钟周期延迟。当R/W依赖或寄存器内存冲突时会有延迟，但是活动线程数>=192时，可以忽略此延迟。

读取模式：读写。

生命周期：生命周期与线程一致。

优化方法：当每个块的线程数为64倍时，获得最佳效果。

2.2 本地内存

修饰符：无

范围：线程私有

速度：慢

读取模式：读写。

当寄存器用完时，将需要存储的值放在本地内存中。

2.3 共享内存

修饰符：__shared__

共享范围：块内所有线程。

速度：快，1个是时钟周期延迟。当组冲突时会有延迟，没有冲突时，会和寄存器一样快。

生命周期：整个线程块。

读取模式：读写。

缺点：有库冲突，半个warp中，多个线程访问相同的库。必须序列化访问才能解决冲突。

2.4 常量内存

修饰符：__constant__，必须在全局空间内和所有核函数之外声明。

共享范围：网格内，可以从主机访问

速度：速度与芯片上的缓存有关；慢速时，1次缓存未命中，所以从全局内存读取；快速时，都命中缓存。

读取模式：只读。

生命周期：随着应用程序结束而结束。

空间：64KB；

初始化：在host端使用 cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src,size_t count);初始化

与常量内存进行数据交换

使用案例：

__constant__ float const_data[256];
float data[256];
cudaMemcpyToSymbol(const_data, data, sizeof(data));
cudaMemcpyFromSymbol(data, const_data, sizeof(data));

2.5 纹理内存

共享范围：网格内，可以从主机访问。

速度与常量内存一致。

读取模式：只读。

生命周期：随着应用程序结束而结束。

块共享内存，块内所有线程都可访问，生命周期与块一致。

全局内存，所有线程都可以访问，生命周期随着应用程序结束而结束。

2.6 全局内存

修饰符：__device__或者在host端使用cudaMalloc函数动态声明。

共享范围：全局。

速度：慢，500个时钟周期。联合访问时会有冲突

读取模式：读写。

生命周期：随着应用程序结束而结束。

初始化：cudaMalloc

3 内存操作函数

3.1 内存拷贝

3.1.1 cudaMemcpyToSymbol——将host拷贝到global内存

cudaMemcpyToSymbol(target,&src,sizeof(src));

若申请的是device内存，cudaMemcpyToSymbol就是从host拷贝到global内存

若申请的是constant内存，cudaMemcpyToSymbol就是从host拷贝到constant内存

复制到device内存测试代码（constant也是如此）：

__device__ int CMTSTest[10];
__global__ void CMTSTestFunction(){
        int i;
        printf("the array is:\n");
        for(i=0;i<10;i++) printf("%d ",CMTSTest[i]);
        printf("\n");
}
void cudaMemcpyToSymbolTest(){
        int *array=(int *)malloc(sizeof(int)*10);
        int i =0;
        for(i=0;i<10;i++) array[i]=i+1;
        cout<<"the target is"<<endl;
        for(i=0;i<10;i++) cout<<array[i]<<" ";
        cout<<endl;

        cout<<"before copy"<<endl;
        CMTSTestFunction<<<1,1>>>();
        cudaDeviceSynchronize();
        cudaError_t err= cudaMemcpyToSymbol(CMTSTest,array,sizeof(int)*10);
        if(err!=cudaSuccess){
                printf("\nError:%s\n",cudaGetErrorString(err));

        }
        cudaDeviceSynchronize();
        cout<<"\nafter copy"<<endl;
        CMTSTestFunction<<<1,1>>>();
        cudaDeviceSynchronize();
        printf("\n");
}

3.1.2 cudaMemcpyFromSymbol——将global内存拷贝到host

cudaMemcpyFromSymbol(&target,src,sizeof(src));

将device变量拷贝到host，字节数为源变量的大小，host端变量传递的是地址。

3.1.3 cudaMemcpy——内存拷贝

cudaMemcpy(void *dst, const void *src, size_t count,TransformType)

TransformType类型有：

- - - cudaMemcpyDeviceToHost：设备到主机
    - cudaMemcpyHostToDevice：主机到设备

特点：同步函数，未完成数据的转移操作之前会一直占有CPU进程控制权。如果想异步调用，使用cudaMemcpyAsync()；

3.2 内存分配

3.2.1 host端内存分配——malloc，cudaMallocHost，free

host内存模式

1. 1. 1. pageable memory：通过操作系统API（malloc()，new()）分配的存储器空间，参与页交换；
    2. pinned memory ：始终存在于物理内存中，不会被分配到低速的虚拟内存中，能够通过DMA加速与设备端进行通信，不参与页交换；
      1. 优点：主机端-设备端的数据传输带宽高，从GPU直接访问，省掉主存与显存间进行数据拷贝的工作；
      2. 缺点：分配过多会导致操作系统用于分页的物理内存变少，导致系统整体性能下降。

malloc分配

直接使用malloc对host端内存进行动态分配，使用free释放指针。

特点：分配的内存是Pageable的（可交换页的）

cudaMallocHost分配

函数：cudaMallocHost((void **)&h_aPinned,int size);

特点：分配的内存是Pinned的（页锁定的）,使用cudaFreeHost()释放

测试代码：

void cudaMallocHostTest(){
        int *pinnedVarabal =NULL;
        cudaMallocHost((int **)&pinnedVarabal,sizeof(int)*10);
        cudaMemset(pinnedVarabal,0,sizeof(int)*10);
        int i =0;
        for(i=0;i<10;i++) pinnedVarabal[i]=i;
         for(i=0;i<10;i++) cout<<pinnedVarabal[i]<<" "; 
}

3.2.2 GPU端内存分配——cudaMalloc，cudaMallocManaged，cudaFree

GPU内存分配——cudaMalloc

函数：cudaError_t cudaMalloc((void **)&devPtr,size_t count)

参数说明：devPtr是指针的指针。

注意：给指针分配了GPU内存后，只能在GPU端使用，在host端调用会报错。（使用cudaMallocManged分配后在host端则不会报错）

测试代码：

int *TestNum=NULL;
size_t TestNumSize = sizeof(int);
cudaMalloc((int**)&TestNum,TestNumSize);
cudaMemset(TestNum,0,TestNumSize);

统一内存分配——cudaMallocManaged

cudaError_t cudaMallocManaged((void **)&devPtr,size_t count)

优点：统一内存可以从系统中的任何处理器访问的单个内存地址空间。分配的内存在host端和device端都可以调用。

缺点：从主机到GPU传输数据略慢与cudaMalloc

测试代码：

__global__ void GPUFunction(int *a){
        printf("A on Gpu is:%d\n",*a);

}
void cudaMallocManagedTest(){
        int *TestNum=NULL;   
        size_t TestNumSize = sizeof(int);
        cudaMallocManaged((int**)&TestNum,TestNumSize);
        cudaMemset(TestNum,0,TestNumSize);
        cout<<"A on CPU is:"<<*TestNum<<endl;
        GPUFunction<<<1,1>>>(TestNum);
}

cudaFree释放GPU内存

cudaError_t cudaFree(void * devPtr)

3.3 内存填充

3.3.1 填充host端内存

使用memset函数即可

3.3.1 填充GPU内存

cudaError_t cudaMemset(void * devPtr, int value ,size_t count)

操作的内存应当是GPU上的内存

3.4 重置关联资源

cudaDeviceReset()

重置当前线程所关联过的当前设备的所有资源

4 优化

使用cudaMallocManaged分配内存时，可以实现自动内存迁移，无需深入了解cudaMallocManaged分配统一内存的工作原理。

4.1 统一内存的迁移

分配统一内存时，内存可能尚未驻留在主机或设备上，主机或设备尝试访问会发生页错误，此时主机或设备会尝试批量迁移数据。CPU和GPU也是如此，当访问未驻留其上的内存时，会发生页错误并触发迁移。

仅通过CPU访问统一内存，或者仅通过GPU访问统一内存的时候，并不会发生内存迁移。但是若先通过CPU访问统一内存，再通过GPU访问该内存，则会发生页错误并且会迁移内存。使用nsys工具分析，发生页错误并迁移内存的时候，则会打印CUDA Memory Operation统计这一栏，名字会显示从哪里迁移至哪里。

4.2 异步内存预取

4.2.1 作用

可以在使用统一内存之前，在后台将其迁移到CPU或者GPU，以此减少页错误以及数据迁移带来的成本。

4.2.2 函数

cudaMemPrefetchAsync(pointerToSomeUMData,size,deviceId); //预取到GPU 
cudaMemPrefetchAsync(pointerToSomeUMData,size,cudaCPUDeviceId);  //预取到CPU

矩阵加法，分配空间后不预取到GPU端耗时：

预取到GPU端耗时：

4.3 块内线程数的选择

块中的线程数必须是warp大小的倍数。（避免填充不足）

最小值：每个块64个线程。

128到258个线程是更好的选择。

4.4 隐藏内存访问延迟

全局内存每次访问需要400-600周期的延迟。

补救方法：

增加块内的线程数。

合并对相邻地址的内存访问。

4次全局访问内存需要花费4*400=1600个周期。

4个并发线程，每个线程1次读取，最少需要400个周期。

4.5 最大化内存吞吐量

全局内存在设备内存中，可通过 32、64 或 128 个字节的规格进行内存访问。这些内存规格必须天然地对齐：只有与其大小对齐的 32、64 或 128 字节（即其第一个地址是其大小的倍数）的设备内存段才能通过内存事务进行读取或写入。

因此，要最大限度地提高全局内存吞吐量，必须通过：

遵循基于计算能力 3.x 、计算能力 5.x、计算能力 6.x、计算能力 7.x 和计算能力 8.x 的最佳访问模式
使用符合下面大小和对齐要求部分中详细说明的大小和对齐要求的数据类型，
在某些情况下，例如，在访问下面的二维矩阵部分中描述的二维矩阵时，应修补数据。

4.6 coalesce

4.6.1 以矩阵为例

情形1

假设矩阵M中，每个线程负责一行。

第0次迭代时，每个线程读取每行的第0个元素。

第1次迭代时，每个线程读取每行的第1个元素。

第2次迭代时，每个线程读取每行的第2个元素。

但是矩阵是按照行优先存储的。

实际上在内存上的顺序是N_0,1,N_0,2,N_0,3,...,N_1,0,N_1,1,N_1,2,...,N_2,0,N_2,1,N_2,2...

这样读取时每个线程之间隔了N-1个元素

情形2

假设矩阵N中，每个线程负责一列。

第0次迭代时，每个线程读取每列的第0个元素。

第1次迭代时，每个线程读取每列的第1个元素。

第2次迭代时，每个线程读取每行的第2个元素。

矩阵是按照行优先存储的。

实际上在内存上的顺序是N_0,1,N_0,2,N_0,3,...,N_1,0,N_1,1,N_1,2,...,N_2,0,N_2,1,N_2,2...

这样相邻的线程就访问了相邻的数据。形成了coalesced合并。

4.6.2 以读取数为例

每一个线程操作一个读取一个float3类型

但是float3类型内部又包含了3个值。所以相邻的线程读取的不是相邻的内存。

读取模式如下：

解决方法1

对于同一个float3变量,使用三个线程分别读取其三个标量保存到共享内存。

每个块有256个线程，线程块也就需要sizeof(float3) x 256个字节的SMEM

将float3的转换成float类型，保存在数组中，也就是输入数组中连续的3个元素为一个float3.

复制到共享内存

计算

从共享内存取出

假设blockDim=256，也就是每个块256个线程

Thread 0: index =3* 0*256 +0 = 0

s_data[0] = g_in[0] , s_data[256] = g_in[256] , s_data[512] = g_in[512]

Thread 1: index = 3* 0*256 +1= 1

s_data[1] = g_in[1] , s_data[257] = g_in[257] , s_data[513] = g_in[513]

Thread 2:

s_data[2] = g_in[2] , s_data[258] = g_in[258] , s_data[514] = g_in[514]

Thread 3:

s_data[3] = g_in[3] , s_data[259] = g_in[259] , s_data[515] = g_in[512]

Thread 4:

s_data[4] = g_in[4] , s_data[260] = g_in[260] , s_data[516] = g_in[516]

Thread 5:

s_data[5] = g_in[5] , s_data[261] = g_in[261] , s_data[517] = g_in[517]

在block0中

Thread 0负责s_data的 0 256 512。和g_in和g_out的0 256 512

...

Thread 255 负责s_data的255 511 767。和g_in和g_out的 255 511 767

在block1中

Thread 0负责s_data的0 256 512。和g_in和g_out的768 1024 1280

....

Thread 255 负责s_data的255 511 767。和g_in和g_out的1023 1279 1535

注意：sharedmemory是块内共享的。

创建一个256*3的__shared__float是每个块中具有256*3大小的float数组。

所以在block1中，threadIdx.x从0开始计算没问题，因为在block1也是一个全新的__shared__float数组。

需要将所有的g_in放入每个块的共享内存，所以g_in在块间增长的步长要*3

解决方法2

使用数组结构体SOA

Struct SOA{ float x[256] , y[256] , z[256]; };

也就是将所有的x连续存放，y连续存放，z连续存放

这样连续的线程就可以连续的读取了。

对于顺序访问模式，但是sizeof(struct)不等于4，8或16，就使用SOA整列。

解决方法3

使用对其说明符alignment specifiers

4.7 矩阵加法程序

长度为10万的向量相加

#include"stdio.h"
#include"ctime"
#include"iostream"
using namespace std;
void VectorAddInCPU(int *A,int *B,int *C,int VectorSize){
	int i;	
	for(i=0;i<VectorSize;i++)
		C[i]=A[i]+B[i];
}
__global__ void VectorAddInGPU(int *A,int *B,int *C,int VectorSize){
	int i=threadIdx.x + blockIdx.x *blockDim.x;
	int stride =blockDim.x * gridDim.x;
	for(;i<VectorSize;i+=stride)
		C[i]=A[i]+B[i];
}

void InitVector(int *A,int n){
	int i;
	for(i=0;i<n;i++)	A[i]=i+1;
}
__global__ void InitVectorOnGPU(int *A,int n){
	int i=threadIdx.x + blockIdx.x *blockDim.x;
	int stride =blockDim.x * gridDim.x;
	for(;i<n;i+=stride)
		A[i]=i;
}

void PrintVector(int *A,int n){
	int i;
	for(i=0;i<n;i++)	cout<<A[i]<<" ";
	cout<<endl;
}
void GPUAddTest(){
	int deviceId;
	cudaGetDevice(&deviceId);
	cudaDeviceProp props;
	cudaGetDeviceProperties(&props,deviceId);
	int SMs= props.multiProcessorCount;
	int WarpSize = props.warpSize;
 	int threadsPerBlock = WarpSize * 10;
	int Blocks = SMs;
	int *A,*B,*C;
	int size=100000;
	size_t vectorSize = sizeof(int)*size;
	cudaMallocManaged((int **)&A,vectorSize);
	cudaMallocManaged((int **)&B,vectorSize);
	cudaMallocManaged((int **)&C,vectorSize);
	InitVector(A,vectorSize);
	InitVector(B,vectorSize);
	cudaMemset(C,0,vectorSize);
	VectorAddInGPU<<<Blocks,threadsPerBlock>>>(A,B,C,size);
	cudaDeviceSynchronize();
}
void GPUAddPreFetchTest(){
	int deviceId;
	cudaGetDevice(&deviceId);
	cudaDeviceProp props;
	cudaGetDeviceProperties(&props,deviceId);
	int SMs= props.multiProcessorCount;
	int WarpSize = props.warpSize;
 	int threadsPerBlock = WarpSize * 10;
	int Blocks = SMs;
	int *A,*B,*C;
	int size=100000;
	size_t vectorSize = sizeof(int)*size;
	cudaMallocManaged((int **)&A,vectorSize);
	cudaMallocManaged((int **)&B,vectorSize);
	cudaMallocManaged((int **)&C,vectorSize);
	cudaMemPrefetchAsync(A,vectorSize,deviceId);
	cudaMemPrefetchAsync(B,vectorSize,deviceId);
	cudaMemPrefetchAsync(B,vectorSize,deviceId);
	InitVectorOnGPU<<<Blocks,threadsPerBlock>>>(A,vectorSize);
	InitVectorOnGPU<<<Blocks,threadsPerBlock>>>(B,vectorSize);
	cudaDeviceSynchronize();
	cudaMemset(C,0,vectorSize);
	VectorAddInGPU<<<Blocks,threadsPerBlock>>>(A,B,C,size);
	cudaDeviceSynchronize();
}
void CPUAddTest(){
	int *A,*B,*C;
	int size=100000;
	A=(int *)malloc(sizeof(int)*size);
	B=(int *)malloc(sizeof(int)*size);
	C=(int *)malloc(sizeof(int)*size);
	InitVector(A,size);
	InitVector(B,size);
	clock_t start= clock();
	VectorAddInCPU(A,B,C,size);
	clock_t end = clock();
	cout<<double(end - start)/ CLOCKS_PER_SEC*1000000000<<" :ns"<<endl;

}
int main(){
	GPUAddPreFetchTest();
	return 0;
}

4.5.1 CPU上计算耗时

4.5.2 GPU上加速，不预取到GPU耗时

4.5.3 GPU上加速，预取到GPU耗时

posted @ 2023-01-30 17:35 Laplace蒜子阅读(1404) 评论(0) 编辑收藏举报

RedNoseBo

CUDA C++ / 内存管理以及优化

内存

1 内存读写速度

2 内存模型

2.1 寄存器

2.2 本地内存

2.3 共享内存

2.4 常量内存

2.5 纹理内存

2.6 全局内存

3 内存操作函数

3.1 内存拷贝

3.1.1 cudaMemcpyToSymbol——将host拷贝到global内存

3.1.2 cudaMemcpyFromSymbol——将global内存拷贝到host

3.1.3 cudaMemcpy——内存拷贝

3.2 内存分配

3.2.1 host端内存分配——malloc，cudaMallocHost，free

3.2.2 GPU端内存分配——cudaMalloc，cudaMallocManaged，cudaFree

3.3 内存填充

3.4 重置关联资源

4 优化

4.1 统一内存的迁移

4.2 异步内存预取

4.2.1 作用

4.2.2 函数

4.3 块内线程数的选择

4.4 隐藏内存访问延迟

4.5 最大化内存吞吐量

4.6 coalesce

4.6.1 以矩阵为例

4.6.2 以读取数为例

4.7 矩阵加法程序

4.5.1 CPU上计算耗时

4.5.2 GPU上加速，不预取到GPU耗时

4.5.3 GPU上加速，预取到GPU耗时

公告