[转] 初步了解PageLocked host memory的mapped memory功能使用

adagio_chen 2014-09-14 08:42:39

加精

http://blog.csdn.net/ee06b109/article/details/6264506

导言：
大家都知道CUDA 中PageLocked memory 相比portable memory 有着多种优势：
1 在有front-side bus的系统中，pagelocked memory 所提供的host 与device之间的数据传送速度，快得多。
2 kernel execution 和 pagelocked memory 与 device memory 间的数据复制可同时进行（具体有待实验）。
3 一些设备（计算能力2.0及以上），pagelocked memory 可以被映射到设备地址空间（mapped memory），从而减少对数据复制的需求，增加程序运行速度，这一项将是本文考察的重点。
当然pageLocked memory 也有缺点，那就是内存如果占用过多将影响程序的运行速度，因为这一块内存被锁定后无法自由分配给其他线程或程序。当然对于主机内存资源足够用的小e来说根本也不成问题（4G）。

实验目的：
了解pagelocked memory 与Mapped 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;

}