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

巧用CUDA中的pinned memory相关推荐

Cuda中Global memory中coalescing例程解释
Global memory是cuda中最常见的存储类型,又叫做Device memory,位于Host主机区域上,它的生命周期是在整个Grid里面,大约具有500个cycle latency.在cud ...
CUDA中grid、block、thread、warp与SM、SP的关系
首先概括一下这几个概念.其中SM(Streaming Multiprocessor)和SP(streaming Processor)是硬件层次的,其中一个SM可以包含多个SP.thread是一个线程, ...
java如何给一个链表定义和传值_如何在CUDA中为Transformer编写一个PyTorch自定义层...
如今,深度学习模型处于持续的演进中,它们正变得庞大而复杂.研究者们通常通过组合现有的 TensorFlow 或 PyTorch 操作符来发现新的架构.然而,有时候,我们可能需要通过自定义的操作符来实现 ...
[原]CUDA中grid、block、thread、warp与SM、SP的关系
[原]CUDA中grid.block.thread.warp与SM.SP的关系 2015-3-27阅读209 评论0 首先概括一下这几个概念.其中SM(Streaming Multiprocessor ...
CUDA中并行规约（Parallel Reduction）的优化
Parallel Reduction是NVIDIA-CUDA自带的例子,也几乎是所有CUDA学习者的的必看算法.在这个算法的优化中,Mark Harris为我们实现了7种不同的优化版本,将Bandwi ...
CUDA中SM对线程块的调度
sm流处理器簇对blocks的调度策略在cuda中,GPU中的SM(比如GTX650有两个SM处理器)被CPU调度器把线程块逐个分配到SM上,每个SM同时处理这个被分配的线程块,但是每次每个时刻只能 ...
CUDA 中 FFT 的使用
CUDA 中 FFT 的使用 @(10.CUDA)[CUDA,并行,fft] 1. 流程使用cufftHandle创建句柄使用cufftPlan1d(),cufftPlan3d(),cufftPl ...
python中接口测试垃圾数据如何清理_巧用PyUnit中unittest特性解决接口测试产生脏数据问题...
巧用PyUnit中unittest特性解决接口测试产生脏数据问题一.背景测试数据创建后需要对其删除,不然可能产生脏数据,对开发和测试.生产环境造成一定影响.其接口框架是基于Python,API规范 ...
CUDA中的一些基本概念
线程线程是CUDA中并行程序的基本构建,一个线程就是程序中国的一个单一的执行流,就像一件衣服上的一块棉,一块块棉交织在一起组成衣服,同样一个个线程组成成并行程序. 随着处理器的核越来越多,硬件可以 ...
CUDA中的复数定义、开内存空间以及运算
最近在做时频混合域的全波形反演(FWI),用CUDA加速,要做复数运算,所以研究了一下CUDA中复数运算等.简单说一下CUDA中复数的基本应用. 在CUDA中用CUFFT的库来定义与运算复数,基本如下 ...

巧用CUDA中的pinned memory

巧用CUDA中的pinned memory相关推荐

最新文章

热门文章