Heterogeneous Parallel Programming(异构并行编程)学习笔记(一)

好记性不如烂笔记。以下是在Coursera学习Heterogeneous Parallel Programming时记录的一些要点。


Wiki对Heterogeneous Programming的解释如下:

Heterogeneous computing systems refer to electronic systems that use a variety of different types of computational units. A computational unit could be a general-purpose processor (GPP), a special-purpose processor (i.e. digital signal processor (DSP) or graphics processing unit (GPU)), a co-processor, or custom acceleration logic (application-specific integrated circuit (ASIC) or field-programmable gate array (FPGA)). 

简要的说,就是采用不同类型的计算节点协同进行计算。而Heterogeneous Parallel Programming则是建立在这种机制上的并行计算。这里使用的是的CUDA。CUDA是NVIDIA推出的建立在C语言和GPU基础上的计算框架。详细情况可参考《NVidia CUDA C Programming Guide》。


1. GPU与CPU

GPU与CPU的设计理念不同:GPU旨在提供高吞吐量,而CPU旨在提供低延迟的操作,如下图所示:


CPU需要降低指令的执行时间,所以有很大的缓存,而GPU则不然。单一的GPU线程执行时间相当长,因此总是多线程并行,这样提高了吞吐量。

综上,在串行计算部分应该使用CPU,而并行计算部分则应使用GPU。


2. CUDA计算模型

CUDA中计算分为两部分,串行部分在Host上执行,即CPU,而并行部分在Device上执行,即GPU。

相比传统的C语言,CUDA增加了一些扩展,包括了库和关键字。CUDA代码提交给NVCC编译器,该编译器将代码分为Host代码和Device代码两部分。Host代码即为原本的C语言,交由GCC或其他的编译器处理;Device代码部分交给一个称为实时(Just in time)编译器的组件,在给代码运行之前编译。


3. Device上的并行线程阵列

并行线程阵列由Grid——Block——Thread三级结构组成,如下图所示:


每一个Grid中包含N个Block,每一个Block中包含N个Thread。

这里需要提到SPMD概念:SPMD,即Single Program Multiple Data,指相同的程序处理不同的数据。在Device端执行的线程即属于此类型,每个Grid中的所有线程执行相同的程序(共享PC和IR指针)。但是这些线程需要从共享的存储中取得自身的数据,这样就需要一种数据定位机制。CUDA的定位公式如下:

i = blockIdx.x * blockDim.x + threadIdx.x

bllockIdx标识Block,blockDim为Block在该维度上的大小,threadIdx为在Block内部线程的标识。

注意到后缀的.x,这是因为CUDA的线程阵列可以是多维的(如上图),blockIdx和threadIdx最多可以达到3维。这能够为处理图像和空间数据提供极大的便利。


4. Device上的内存模型

Device上的内存模型如下图所示:


每个Grid有一个共享的存储,其中每个线程有自己的寄存器。Host代码负责分配Grid中的共享内存空间,以及数据在Host、Device之间的传输。Device代码则只与共享内存、本地寄存器交互。


5. CUDA基本函数

与C语言中的函数想对应,CUDA有以下几个基本函数:

cudaMalloc()、cudaFree()、cudaMemcpy()

其作用等同于C中的对应函数,不同之处在于这些函数操作的是Device中的共享内存,cudaMemcpy()则用于Host内存与Device内存传输数据。


6. 函数标识

CUDA的函数分为三种:


注意都是双下划线。其中的__global__函数即为C代码中调用Device上计算的入口。

__host__函数为传统的C函数,也是默认的函数类型。之所以增加这一标识的原因是有时候可能__device__和__host__共同使用,这时可以让编译器知道,需要编译两个版本的函数。


7. 例:向量加法

注意:这不是一个完整的程序,只用于体现主要步骤

#include<cuda.h>

...

// vector addition

void vecAdd(float *h_A, float *h_B, float *h_C, int n)

{

int size = n * sizeof(float);

float *d_A;

float *d_B;

float *d_C;


// allocate device memory for A, B, C

cudaMalloc((void**)&d_A, size);

cudaMalloc((void**)&d_B, size);

cudaMalloc((void**)&d_C, size);


// transfer data from host to device

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);


// specify the dimension for Grid and Block first

// 256 threads in one block, and make sure there are enough grid to cover all the vector

dim3 gridDim((n-1) / 256, 1, 1);

dim3 blockDim(256, 1, 1);


// kernel invocation code

vecAddKernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, n);


// transfer result from device to host

cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// free memory

cudaFree(d_A);

cudaFree(d_B);

cudaFree(d_C);

}


// vector addition kernel function with __global__ identifier

__global__

void vecAddKernel(float *d_A, float *d_B, float *d_C, int n)

{

// calculate index

int i = blockIdx.x * blockDim.x + threadIdx.x;

if (i < n)

d_C[i] = d_A[i] + d_B[i];

}


阅读更多
文章标签: cuda parallel 编程
想对作者说点什么? 我来说一句

没有更多推荐了,返回首页

关闭
关闭
关闭