CUDA C编程7： 内存管理之固定内存

系列文章目录

---

---

前言

这里开始介绍CUDA C编程内存管理中的固定内存相关知识。

---

一、固定内存相关知识点

系统分配的主机内存默认是可分页的，GPU不能在可分页主机内存上安全地访问数据，这是因为主机操作系统在物理位置上移动数据时，GPU无法控制。

当从可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存。

使用如下函数可直接分配固定主机内存：
$cudaError\_tcudaMallocHost(void\ast \ast devPtr,size\_tcount);$

固定内存是页面锁定的并且对设备来说是可访问的，由于固定内存能被设备直接访问，所以允许更好的带宽进行读写。

需要注意的是，分配过多的固定内存可能会降低主机系统的性能，这是因为它减少了用于存储虚拟内存数据的可分页内存的数量。

使用如下函数释放固定主机内存：
$cudaError\_tcudaFreeHost(void\ast ptr);$

CPU与GPU通过固定内存传输数据流程如下图所示（引用自参考资料）：

二、固定内存示例

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

int main()
{
	//set up device
	int dev = 0;
	cudaSetDevice(dev);

	//memory size
	unsigned int isize = 1 << 22;
	unsigned int nbytes = isize * sizeof(float);

	//get device information
	cudaDeviceProp stDeviceProp;
	cudaGetDeviceProperties(&stDeviceProp, dev);
	printf("starting at ");
	printf("device %d: %s memory size %d nbyte %5.2fMB\n", dev,
		stDeviceProp.name, isize, nbytes / (1024.0f * 1024.0f));

	//allocate the host memory
	//float* h_a = (float*)malloc(nbytes);
	float* h_a = nullptr;
	cudaError_t err = cudaMallocHost(&h_a, nbytes);//fixed page global memory
	if (err != cudaSuccess)
	{
		printf("cudaMallocHost failed!\n");
		goto EXIT;
	}

	//allocate the device memory
	float* d_a;
	cudaMalloc(&d_a, nbytes);//pageable global memory

	//initialize the host memory
	for (unsigned int i = 0; i < isize; i++)
		h_a[i] = 0.5f;

	//transfer data from the host to the device
	cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);
	//cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToHost);

	//tansfer data from the device to the host
	cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);
	//cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyHostToHost);

	EXIT:
	//free memory
	
	if (d_a != nullptr)
	{
		cudaFree(d_a);
	}

	if (h_a != nullptr)
	{
		cudaFreeHost(h_a);
		//free(h_a);
	}
		

	//reset device
	cudaDeviceReset();

	system("pause");
	return 0;
}

可分页的全局内存程序nvprof分析结果：
固定内存程序nvprof分析结果：

对比固定内存与可分页内存运行时间可知，固定内存访问速度(Duration)更快，吞吐量（Throughput）更大。

---

总结

使用固定内存在CPU与GPU间传输数据有更高的性能，但是要慎用。本人也是现学现卖，如果有任何失误之处，欢迎大神们指正。

参考资料

《CUDA C编程权威指南》