《CUDA编程》6.CUDA的内存组织

最新推荐文章于 2025-07-11 11:07:36 发布

原创

最新推荐文章于 2025-07-11 11:07:36 发布 · 1.2k 阅读

24 ·

CC 4.0 BY-SA版权

文章标签：

#c++ #cuda

前面几章讲了一些编写高性能CUDA程序的要点，但还有很多其他需要注意的，其中最重要的就是合理的使用设备内存

1 CUDA的内存组织简介

现代计算机中的内存存在一种组织结构(hierachy)，即不同类型的内存具有不同的容量和访问延迟（可以理解为处理器等待内存的时间）。一般来说，延迟低的内存容量小，延迟高的内存容量大。

下表是CUDA设备（显卡）中的几种内存和主要特征
在这里插入图片描述
下图是组织示意图和数据移动发方向图

2 不同内存的简介

2.1 全局内存

全局内存（global memory）： 核函数中的所有线程都能够访问。

特点：

容量大： 通常是GPU上最大的内存区域。
访问延迟高： 相比于其他内存，访问延迟较高。
全局可见： 对所有线程和线程块可见，可在核函数中自由读取、可读可写。
持久性： 全局内存中的数据在内核函数执行期间保持不变，可以跨多个内核调用使用。

全局内存主要是为核函数提供数据，并在主机和设备、设备和设备之间传输数据

全局内存的生命周期由主机端决定，所以：

cudaMalloc()函数是主机在设备的全局内存中分配一段指定大小的内存区域
cudaFree()函数是主机把该内存释放

以上所说的全局内存称为线性内存（linear memory），还有一种不对用户透明的内存称为CUDA Array，专为纹理拾取服务。

和C++函数一样，cudaMalloc()函数是动态地分配内存，CUDA中也允许使用静态全局内存变量，定义方法如下：

__device__ T x; 定义单个变量
__device__ T y[N]; 定义固定长度地数组

其中，修饰符 __ device __ 说明该变量是设备中的变量，而不是主机中的变量；T 是变量的类型；

在核函数中，可以直接访问静态全局内存变量，不需要以参数的形式传入；注意，主机函数无法直接访问，只能用cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()在主机内存和静态全局内存之间传输数据，下面是两个函数的结构：

①将主机数据复制到静态全局内存中
在这里插入图片描述
②将静态全局内存数据复制到主机中

之后会讨论一种利用静态全局内存加速程序的技巧，现在给出使用例子：

#include <cuda.h>
#include <cuda_runtime.h>
#include "error_check.cuh"

__device__ int d_x = 1;
__device__ int d_y[2];

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

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

	cudaOut<<<1,1>>>();
	CHECK(cudaDeviceSynchronize());

	CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));
	printf("Host:  h_y[0]=%d, h_y[1]=%d\n", h_y[0], h_y[1]);

	return 0;
}