【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;
}


评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值