系列文章目录

文章目录

系列文章目录
前言
一、零拷贝内存相关知识点
二、零拷贝内存示例
- 1. 代码实现
- 2. 运行结果
总结
参考资料

前言

这里跟大家分享内存管理第三篇：零拷贝内存。

一、零拷贝内存相关知识点

之前学习的CUDA知识中，主机不能直接访问设备变量，需要通过 $c u d a M e m c p y$ 函数实现主机与设备间数据拷贝，当然设备也不能直接访问主机变量。

这里介绍的零拷贝内存则是个例外，主机和设备都可以访问零拷贝内存。

注意，零拷贝内存相当于从全局内存中分出的一块独立内存，使用了固定内存技术实现零内存拷贝。

在CUDA核函数中使用零拷贝内存的优势如下：
（1）当设备内存不足时可利用主机内存
（2）避免主机和设备间的显示数据传输
（3）提高PCIe传输率

因为设备和主机都可以访问零拷贝内存数据，那么就要注意同步问题，避免主机和设备同时更改零拷贝内存中的数据，否则将产生脏数据。

零拷贝内存的技术实现主要靠固定内存（不可分页），该内存映射到设备地址空间中。可通过如下函数创建一个到固定内存的映射：
$cudaError\_t\ cudaHostAlloc(void\ **pHost,\ size\_t\ count,\ unsigned\ int\ flags);$

零拷贝内存需要用 $c u d a F r e e H o s t$ 函数释放。

分配零拷贝内存的flags列举如下：
（1）cudaHostAllocDefault
cudaHostAlloc函数行为与分配固定内存函数cudaMallocHost函数一致；
（2）cudaHostAllocPortable
使函数返回能被所有CUDA上下文使用的固定内存，而不仅是执行内存分配的那一个。
（3）cudaHostAllocWriteCombined
使函数返回写结合内存，该内存可在某些系统配置（哪些系统配置呢？）通过PCIe总线上更快地传输，但是它在大多数主机上不能被有效地读取。
（4）cudaHostAllocMapped
这是零拷贝内存的最明显标志，可以实现主机写入和设备读取被映射到设备地址空间中的主机内存。

可通过下列函数获取映射到固定内存的设备指针：
$cudaError\_t cudaHostGetDevicePointer(void\ **pDevice,\ void\ *pHost,\ unsigned\ int\ flags);$

上述函数返回设备指针，该指针可在设备上被引用以访问映射得到的固定主机内存。

注意，如果设备不支持映射得到的固定内存，上述函数将失效。

在进行频繁的读写操作时，使用零拷贝内存作为设备内存的补充将显著降低性能。因为每一次映射到内存的传输必须经过PCIe总线。与全局内存相比，延迟也显著增加。

二、零拷贝内存示例

1. 代码实现

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

void InitData(float* data, size_t nElem)
{
	for (size_t i = 0; i < nElem; i++)
	{
		data[i] = i % 255;
	}
}

void SumArraysOnHost(float* h_A, float* h_B, float* hostRef, size_t nElem)
{
	for (size_t i = 0; i < nElem; i++)
	{
		hostRef[i] = h_A[i] + h_B[i];
	}
}

void CheckResults(float* hostRef, float* gpuRef, size_t nElem)
{
	bool bSame = true;
	for (size_t i = 0; i < nElem; i++)
	{
		if (abs(gpuRef[i] - hostRef[i]) > 1e-5)
		{
			bSame = false;
		}
	}

	if (bSame)
	{
		printf("Result is correct!\n");
	}
	else
	{
		printf("Result is error!\n");
	}
}

__global__ void GpuSumArrays(float* d_A, float* d_B, float* d_C, size_t nElem)
{
	int tid = blockDim.x * blockIdx.x + threadIdx.x;
	if (tid < nElem)
		d_C[tid] = d_A[tid] + d_B[tid];
}




int main()
{
	int nDev = 0;
	cudaSetDevice(nDev);

	cudaDeviceProp stDeviceProp;
	cudaGetDeviceProperties(&stDeviceProp, nDev);

	//check whether support mapped memory
	if (!stDeviceProp.canMapHostMemory)
	{
		printf("Device %d does not support mapping CPU host memory!\n", nDev);
		goto EXIT;
	}

	printf("Using device %d: %s\n", nDev, stDeviceProp.name);

	// set up data size of vector
	int nPower = 10;
	int nElem = 1 << nPower;
	size_t nBytes = nElem * sizeof(float);
	if (nPower < 18) {
		printf("Vector size %d power %d nbytes %3.0f KB\n",
			nElem, nPower, (float)nBytes / (1024.0f));
	}
	else {
		printf("Vector size %d power %d nbytes %3.0f MB\n",
			nElem, nPower, (float)nBytes / (1024.0f * 1024.0f));
	}

	// part 1: using device memory
	// malloc host memory
	float *h_A, *h_B, *hostRef, *gpuRef;
	h_A = (float*)malloc(nBytes);
	h_B = (float*)malloc(nBytes);
	hostRef = (float*)malloc(nBytes);
	gpuRef = (float*)malloc(nBytes);

	// initialize data at host side
	InitData(h_A, nElem);
	InitData(h_B, nElem);
	memset(hostRef, 0, nBytes);
	memset(gpuRef, 0, nBytes);

	// add vector at host side for result checks
	SumArraysOnHost(h_A, h_B, hostRef, nElem);

	// malloc device global memory
	float* d_A, *d_B, *d_C;
	cudaMalloc(&d_A, nBytes);
	cudaMalloc(&d_B, nBytes);
	cudaMalloc(&d_C, nBytes);

	//transfer data from host to device
	cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

	// set up execution configuration
	int nLen = 512;
	dim3 block(nLen);
	dim3 grid((nElem + block.x - 1) / block.x);

	//invoke kernel at host side
	GpuSumArrays << <grid, block >> > (d_A, d_B, d_C, nElem);

	//copy kernel result back to host side
	cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

	//check device results
	CheckResults(hostRef, gpuRef, nElem);

	// free device globl memory
	cudaFree(d_A);
	cudaFree(d_B);
	free(h_A);
	free(h_B);

	// part2: using zerocopy memory for array A and B
	// allocate zerocpy memory
	unsigned int nFlags = cudaHostAllocMapped;
	cudaHostAlloc(&h_A, nBytes, nFlags);
	cudaHostAlloc(&h_B, nBytes, nFlags);

	// initialize data at host side
	InitData(h_A, nElem);
	InitData(h_B, nElem);
	memset(hostRef, 0, nBytes);
	memset(gpuRef, 0, nBytes);

	// pass the pointer to device
	cudaHostGetDevicePointer(&d_A, h_A, 0);
	cudaHostGetDevicePointer(&d_B, h_B, 0);

	// add at host side for result checks
	SumArraysOnHost(h_A, h_B, hostRef, nElem);

	//execute kernle with zero copy memory
	GpuSumArrays << <grid, block >> > (d_A, d_B, d_C, nElem);

	//copy kernel result back to host side
	cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

	//check device results
	CheckResults(hostRef, gpuRef, nElem);

	// free memory
	cudaFree(d_C);
	cudaFreeHost(h_A);
	cudaFreeHost(h_B);

	free(hostRef);
	free(gpuRef);

EXIT:
	cudaDeviceReset();

	system("pause");
	return 0;
}