【CUDA】初步了解PageLocked host memory的mapped memory功能使用

导言：

    大家都知道CUDA 中PageLocked memory 相比portable memory 有着多种优势：

1. 在有front-side bus的系统中，pagelocked memory 所提供的host 与device之间的数据传送速度，快得多。测试结果如图Fig.1 Fig.2所示。
2. kernel execution 和 pagelocked memory 与 device memory 间的数据复制可同时进行（具体有待实验）。
3. 一些设备（计算能力2.0及以上），pagelocked memory 可以被映射到设备地址空间（mapped memory），从而减少对数据复制的需求，增加程序运行速度，这一项将是本文考察的重点。

    当然pageLocked memory 也有缺点，那就是内存如果占用过多将影响程序的运行速度，因为这一块内存被锁定后无法自由分配给其他线程或程序。当然对于主机内存资源足够用的小e来说根本也不成问题（4G）。

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

Fig.1&2 测试结果来源于SDK中bandwidthTest程序

 

实验目的：

    了解pagelocked memory 与Mapped memory的用法。

 

实验过程： 

    1、书写代码：建立基于Mapped pagelocked memory 的应用程序需要以下几步

    第一步：

    分别创建希望通过GPU处理的数据的host端指针，与设备端指针，在下面程序中为：int* h_pData 与 int* d_pData.

    第二步：

    通过cudaSetDeviceFlags置"cudaDeviceMapHost"标志。

    第三部：

    通过cudaHostAlloc函数分配pagelocked memory 给指针“h_pData”。

    第四部：

    通过cudaHostGetDevicePointer函数将主机内存映射到设备端，并绑定指针“d_pData”。

    第五部：

    执行kernel程序（注意用d_pData）

    第六部：

    因为设备主机共用pagelocked memory主机端不用内存传输就可直接显示结果。但也正因为不用内存传输（部分内存传输具有同步功能如：cudaMemcpy，部分异步如：cudaMemoryAsync），kernel会在没运行完之前就返回控制给host,所以要想正常显示结果就要使用cudaThreadSynchronize()函数加以控制。

#include <stdlib.h>
#include <stdio.h>
#include <cutil_inline.h>
#include <cuda.h>
#include <shrUtils.h>
#include <assert.h>  


__global__ void cu_arrayDelete(int* arrayIO)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	arrayIO[idx] = arrayIO[idx] - 16;
}
void checkCUDAError(const char *msg)   
{   
  cudaError_t err = cudaGetLastError();   
  if( cudaSuccess != err) {   
    printf("Cuda error: %s: %s./n", msg, cudaGetErrorString( err) );   
    exit(EXIT_FAILURE);   
  }                            
} 
int main(int argc, char *argv[])
{
	int* h_pData;
	int* d_pData;
	cudaDeviceProp deviceProp;  
	cudaGetDeviceProperties(&deviceProp, 0);     
	if(!deviceProp.canMapHostMemory) {   
		printf("Device %d cannot map host memory!/n");      
	}   
	cudaSetDeviceFlags(cudaDeviceMapHost);
	checkCUDAError("cudaSetDeviceFlags");   

	cutilSafeCall(cudaHostAlloc((void**)&h_pData, 512, cudaHostAllocMapped));
	cudaHostGetDevicePointer((void **)&d_pData, (void *)h_pData, 0);
	for(int i=0; i<128; i++)
	{
		h_pData[i] = 255;
	}
	cu_arrayDelete<<<4,32>>>(d_pData);
	cudaThreadSynchronize();   
	for(int i = 0 ; i<128; i++ )
		printf("%d/n",h_pData[0]);
	while(1);
	return 0;
}