ECE408 Lecture 2 Introduction to Parallel Computing and CUDA

这次回顾ECE408 Lecture 2，这次的内容是并行计算核CUDA的介绍。

课程主页：

https://wiki.illinois.edu/wiki/display/ECE408

搬运视频：

https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz

本讲内容

数据并行计算的基本概念；
CUDA C/C++ 编程接口的基本特性；

数据并行计算示例：将彩色图像转换为灰度图像

CUDA执行模型

典型的计算密集型C/C++ 代码
- 串行或适度并行的部分；
- 高度并行的部分；
串行部件->CPU（或Host）
高度并行部分->GPU（或Device）

并行线程数组

CUDA内核作为线程网格（grid/数组）执行

网格中的所有线程都运行相同的内核代码；
单程序多数据（SPMD 模型）；
每个线程都有一个唯一的索引，用于计算内存地址和做出控制决策；

CUDA的逻辑执行模型

每个CUDA 内核

由grid执行；
grid是thread blocks的3D数组；
thread blocks是threads的3D数组；

示意图：

基本概念

`gridDim`给出block的数量

上图中：

gridDim.x = 8
gridDim.y = 3
gridDim.z = 2

对于1维/2维grid，只要取y, z为1即可。

`blockIdx`对每个block是唯一的

每个block有唯一的索引数组：

blockIdx.x取值0~gridDim.x-1；
blockIdx.y取值0~gridDim.y-1；
blockIdx.z取值0~gridDim.z-1；

`blockDim`是每个block的线程数量

上图中：

blockDim.x = 5
blockDim.y = 4
blockDim.z = 3

对于1维/2维block，只要取y, z为1即可。

`threadIdx`对每个线程唯一

每个线程有唯一的索引数组：

threadIdx.x取值0~blockDim.x-1；
threadIdx.y取值0~blockDim.y-1；
threadIdx.z取值0~blockDim.z-1；

小结

threadIdx元组对于block内的每个线程都是唯一的；
threadIdx和blockIdx对于grid内的每个线程都是唯一的；

线程块：可扩展的协作

将线程数组分成多个block

block内的线程通过共享内存、原子操作和barrier synchronization进行协作（稍后介绍）
不同块中的线程合作较少（稍后介绍）；

图示：

`blockIdx`和`threadIdx`

每个线程使用索引来决定要处理的数据
- blockIdx：1D、2D或3D；
- threadIdx：1D、2D或3D；
在处理多维数据时简化内存寻址
- 图像处理；
- 向量、矩阵、张量；
- 解决体积上的偏微分方程；

示意图：

向量加法

从向量加法了解并行计算。

概念视角

传统的C/C++代码

// Compute vector sum C = A + B
void vecAdd(float* A, float* B, float* C, int n)
{
	for (i = 0, i < n, i++)
		C[i] = A[i] + B[i];
}

int main()
{
    // Memory allocation for A_h, B_h, and C_h
    // I/O to read A_h and B_h, N elements
    …
    vecAdd(A_h, B_h, C_h, N);
}

CUDA `vecAdd` Host代码

#include <cuda.h>

void vecAdd(float* A, float* B, float* C, int n)
{
    int size = n* sizeof(float);
    float* A_d, B_d, C_d;
    
    1. // Allocate device memory for A, B, and C
       // copy A and B to device memory
        
    2. // Kernel launch code – to have the device
       // to perform the actual vector addition
        
    3. // copy C from the device memory
       // Free device vectors
}

CUDA内存的部分概述

Device代码可以
- R/W每个线程寄存器；
- R/W每个网格全局内存；
Host代码可以
- 将数据传入/传出每个网格全局内存；

图示：

CUDA API

cudaMalloc()
- 在device全局内存中分配对象；
- 两个参数；
  - 指向已分配对象的指针地址；
  - 已分配对象的大小（以字节为单位）；
cudaFree()
- 从设备全局内存中释放对象；
- 指向已释放对象的指针；
cudaMemcpy()
- 内存数据传输
- 需要四个参数
  - 指向目的地的指针；
  - 指向源的指针；
  - 复制的字节数；
  - 传输类型/方向；

CUDA `vecAdd` Host代码（续）

void vecAdd(float* A, float* B, float* C, int n)
{
    int size = n * sizeof(float);
    float* A_d, B_d, C_d;
    1. // Transfer A and B to device memory
       cudaMalloc((void **) &A_d, size);
       cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
       cudaMalloc((void **) &B_d, size);
       cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);
       // Allocate device memory for
       cudaMalloc((void **) &C_d, size);
    2. // Kernel invocation code – to be shown later
    3. // Transfer C from device to host
       cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
       // Free device memory for A, B, C
       cudaFree(A_d); cudaFree(B_d); cudaFree (C_d);
}

CUDA `vecAdd` 完整代码

// Compute vector sum C = A+B
// Each thread performs one pair-wise addition
// kernel code
__global__
void vecAddKernel(float* A_d, float* B_d, float* C_d, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
    // 为了防止数组越界，增加判断
	if(i<n) C_d[i] = A_d[i] + B_d[i];
}

// host code
int vectAdd(float* A, float* B, float* C, int n)
{
	// A_d, B_d, C_d allocations and copies omitted
	// Run ceil(n/256) blocks of 256 threads each
    dim3 DimGrid(ceil(n/256), 1, 1);
    dim3 DimBlock(256, 1, 1);
    
	vecAddKernel<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d, n);
}

说明：

DimGrid：网格中每个维度的块数；
DimBlock：块中每个维度的线程数；
blockIdx.x：x维度的块索引；
blockDim.x：x维度中每个块的线程数；
threadIdx.x：块中x维度的线程索引；

整体视角

有关CUDA函数声明的更多信息

	执行在：	只能被调用的位置：
`__device__ float DeviceFunc()`	device	device
`__global__ void KernelFunc()`	device	host
`__host__ float HostFunc()`	host	host

说明：

__global__定义内核函数
- 每个“__”由两个下划线字符组成；
- 内核函数必须返回void；
__device__和__host__可以一起使用；