CUDA并行模型

摘要

本文主要讲讲CUDA并行模型，主要包括：并发线程模型，并行函数调度模型，以及并行内存模型三个部分。

1. 并发线程模型

CUDA的threadIdx排列方式（或者说是block的维度）是空间三维的，而且，uniqueThreadIdx都是唯一编号的，因此不管线程的排列方式与其唯一的编号二者之间存在一个索引的映射方程。

对于一维grid(x维度)，一维block(x维度)的线程排列方式而言，uniqueThreadIdx = blockIdx.x *blockDim.x + threadIdx.x;；

对于一维grid(x维度)，二维(x,y维度)的线程排列方式而言，uniqueThreadIdx =  blockIdx.x * blockDim.x * blockDim.y  
        + threadIdx.y * blockDim.x + threadIdx.x；

更多的维度之间的映射方程见以前的博客：唯一标识threadIdx的索引方程

2. 并行函数调用模型

向量A和向量B加和放到向量C中，该功能CUDA实现函数如下所示：

// Kernel definition __global__ void VecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { ... // Kernel invocation with N threads VecAdd<<<1, N>>>(A, B, C); ... } 

这段代码对于理解并行函数调用模型有着重要的意义！

代码的意思是：任何一个thread(threadIdx=0, 1,2,3, ..., N-1) 都有其独立的计算单元和独立的内存（registers）,当然还有其他的线程之间共享的内存等。换句话说，任何运行时刻，任何并行运行中的某个线程，都在同步独立地执行VecAdd()函数，当然不太可能同一时刻结束，所以以后涉及到线程分歧等问题，这是后话。这是CUDA并行思想的基本实现单元，没看懂一定要反复思考，反复思考，反复思考......直到彻悟。

3. 并行内存模型

如上图所示，说的是每一个thread都有其独立的（别的thread不能碰的！）local memory（即registers），依次类推，每一个block有其独立的shared memory，如图所示，该block中的线程可以访问该shared memory，不同的grid可以访问同一个global memory等等。这是基本的入门内容，所以没看懂一定要反复思考，反复思考，反复思考......直到彻悟。