巧用CUDA中的pinned memory

      这几天看别人的论文,发现一个比较有意思的实现方式。巧用pinned memory,在GPU中实现类似pipeline的功能。在论文中pipeline中,有四个操作:地址生成,数据组装,数据拷贝和计算。对于地址生成和计算是在GPU中操作的。

      详细的请看一个例子:

     1、我们假设有两个thread block,对于第一个block计算地址空间(在例子中省略了),在第一个block生成地址完成生成一个信号; 

     2、当第一个block完成功能后,通知cpu端,此时在cpu端组织对应的数据;

     3、第二部完成后,把数据拷贝到GPU段,数据拷贝完成后给GPU一个信号

    4、GPU中第二个block根据第3部中的信号来计算

    当从论文中看到实现方式时,觉得so easy,然后码代码,结果却不对。研究了两天,并且和论文作者沟通了一番才真正码代码实现了上述功能。现在就上代码来。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

__global__ void pipeline(int *flag_a,int*flag_b,int*Input,int*Out)
{
	int idx=threadIdx.x;
	if(blockIdx.x==0){
		if(0==idx)
			flag_a[0]=1;	//地址生成ok信号
	}
	
	if(blockIdx.x==1){
	
		if(0==idx){
			int value = 0;
		   	do  
		    	{   
		            	asm volatile("ld.global.cg.u32 %0, [%1];" :"=r"(value) :"l"(&flag_b[0]));//数据发送ok信号
		    	} while(value != 1);  
		}
		__syncthreads();
		Out[idx]=Input[idx]+idx;
	}


}
int main()
{
	/*1*/
	int *flag_a,*flag_b;
	cudaHostAlloc((void**)&flag_a,sizeof(int),cudaHostAllocMapped);
	cudaHostAlloc((void**)&flag_b,sizeof(int),cudaHostAllocMapped);
	flag_a[0]=0;
	flag_b[0]=0;
	/*2*/
	int*Input,*Out;
	int *d_Input,*d_Out;
	int*d_float_a,*d_float_b;
	Input=(int*)malloc(sizeof(int)*32);
	Out=(int*)malloc(sizeof(int)*32);
	for(int i=0;i<32;i++){
		Input[i]=i;
	}
	memset(Out,0,sizeof(int)*32);

	cudaMalloc((void**)&d_Input,sizeof(int)*32);
	cudaMemset(d_Input,0,sizeof(int)*32);
	cudaMalloc((void**)&d_Out,sizeof(int)*32);
	cudaMemset(d_Out,0,sizeof(int)*32);

	cudaHostGetDevicePointer((void **)&d_float_a, (void *)flag_a, 0);
	cudaHostGetDevicePointer((void **)&d_float_b, (void *)flag_b, 0);

	cudaStream_t stream_kernel,stream_datacopy;
	cudaStreamCreate(&stream_kernel);
	cudaStreamCreate(&stream_datacopy);

	pipeline<<<2,32,0,stream_kernel>>>(d_float_a,d_float_b,d_Input,d_Out);
	while(!(1==flag_a[0])){	
			cudaMemcpyAsync(d_Input,Input,sizeof(int)*32,cudaMemcpyHostToDevice,stream_datacopy);
			cudaStreamSynchronize(stream_datacopy);
			flag_b[0]=1;
			break;
		
	}

	cudaStreamSynchronize(stream_kernel);
	cudaMemcpy(Out,d_Out,sizeof(int)*32,cudaMemcpyDeviceToHost);
	for(int i=0;i<32;i++){
		printf("%d:%d\n",i,Out[i]);
	}


	cudaFreeHost(flag_a);
	cudaFreeHost(flag_b);
	
	cudaFree(d_Input);
	cudaFree(d_Out);
	free(Out);
	free(Input);
	return 0;

}

  运行结果:

lucas@lucas-desktop:~/cuda/new$ ls
Makefile  pipeline.cu
lucas@lucas-desktop:~/cuda/new$ make
nvcc -O3 -o pile  pipeline.cu
lucas@lucas-desktop:~/cuda/new$ ./pile 
0:0
1:2
2:4
3:6
4:8
5:10
6:12
7:14
8:16
9:18
10:20
11:22
12:24
13:26
14:28
15:30
16:32
17:34
18:36
19:38
20:40
21:42
22:44
23:46
24:48
25:50
26:52
27:54
28:56
29:58
30:60
31:62
lucas@lucas-desktop:~/cuda/new$ 
        由于CPU和GPU端常规方式在kernel执行期间不能直接通信,所以利用了pinned memory,并且利用了两个stream。代码也是比较简单,就不详细说明了。

       需要注意的是在kernel段,为什么会有:

asm volatile("ld.global.cg.u32 %0, [%1];" :"=r"(value) :"l"(&flag_b[0]));//数据发送ok信号

       ptx的代码呢?

      在刚开始的实现中,我的办法如下: 

while(!(1==flag_b[0])){
}	
    对于此种实现方式,在开始实现while之前,编译器优化掉flag_b[0]的值,把其值保存在一个寄存器中。使得pinned memory中flag_b有更新,但是kernel中的值并不是最新的,所以才有ptx的代码。

 



评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值