巧用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的代码。