CUDA C编程9：内存管理之统一虚拟寻址、统一内存寻址

系列文章目录

---

---

前言

这里开始介绍内存管理中的统一虚拟寻址和统一内存寻址技术的相关知识点。

---

一、统一虚拟寻址

1. 统一虚拟寻址技术

只有计算力在2.0及以上版本的设备支持一种特殊的寻址方式，即为统一虚拟寻址（UVA）。有了UVA，主机内存和设备内存共享统一虚拟地址空间，说白了，和之前介绍的零拷贝内存技术的功能相同，不需要$cudaMemcpy$函数完成主机内存与设备内存数据的相互传输，即可以在主机和设备中都可以直接读写。
通过UVA, 有$cudaHostAlloc$分配的固定主机内存具有相同的主机和设备指针。

与零拷贝内存相比， 使用UVA无须获取设备指针或管理物理上数据完全相同的两个指针。从$cudaHostAlloc$函数返回的指针可以直接传递给核函数，这样减少了代码量，提高了应用程序的可读性和可维护性。

2. 示例程序

#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