CUDA编程的框架-以向量相加为例

在这里插入图片描述

CPU适合控制GPU程序的逻辑结构。
注意在编程时要区分CPU的程序和GPU的程序，CPU的内存和GPU的内存。
host- CPU
device- GPU
CPU的内存和GPU的内存之间是相互独立的，因此需要进行通信。

在这里插入图片描述

__global__ //核函数的声明符号<<<grid,block>>>
//在调用核函数时要对核函数进行配置，需要定义它的网格和块的维度（是dim3的数据类型）//在运算时，block中有很多内置的变量，
//比如线程的ID，或者block的ID，
//可以直接调用来获得block或者thread的编号，
//编号可以是1D/2D/3D的（1/2/3维的）

在这里插入图片描述

int main(){int N = 1000000;      //向量的长度是1Mint bs = 256;         //每个块有256个线程int gs = (N+bs-1)/bs; //gs是一共有多少个块，即grid的数目//kernel, call GPUvec_add<<<gs,bs>>>(x,y,z,N);//这里的示例是一维的网格、一维的进程块return 0;
}//kernel定义，必须用__global__标识，返回类型必须是void
__global__ void vec_add(double *x, double *y, double *z, int n){int i = get_tid(); //user-defined macro/function//这里的get_tid()是用户定义的一个宏，调用它可以得到当前线程一维的IDif(i<n) z[i] = x[i]+y[i];
}

GPU上很多并行化是轻量级的线程。GPU的并行是粒度很细的，每个线程可以处理一个数据或者几个数据。

在这里插入图片描述
从程序上（程序员上）分为两级block和grid。
在真正执行时是warp/block/grid三级。

在这里插入图片描述

grid和block都是定义为dim3的类型，我们可以理解为三维的无符号整数。

用dim3的类型构造核函数的示例
在这里插入图片描述
可以根据向量的长度自己计算需要多少网格比较合适，向量很长，网格多一些比较合适。

⬆️因为2）：CPU和GPU运行是异步的，CPU一运行完kernel函数之后就会走，不管GPU有没有算完，所以在写程序的时候要知道GPU算完没。
3）_ _ device _ 说明这个函数只能在GPU中执行，只能在GPU中调用，不能像刚刚一样<<<gs,bs>>>的调用。
必须是GPU的函数调用 _ _ device _ _的函数。
4） _ host _ _的声明一般省略不写。

在这里插入图片描述

CUDA有一些不用申明就能使用的内置变量。

对于block来说，
它需要知道自己的块的ID：blockIdx.x, blockIdx.y, blockIdx.z ，
它需要通过gridDim知道grid在x，y，z方向各有多少个。

对于thread来说，
它需要知道自己的块的ID：threadIdx.x, threadIdx.y, threadIdx.z ，
它需要通过blockDim知道block在x，y，z方向各有多少个。

一个程序块上的线程是放在同一个流多处理器（SM）上的。不能跨流多处理器的放置。

在这里插入图片描述

// 假设block是一维的，grid是二维的，
// 通过定义下面的两个宏可以计算出线程的全局编号和块的全局编号/* get thread id：1D block ans 2D grid */
#define get_tid() (blockDim.x*(blockIdx.x+blockIdx.y*gridDim.x)+threadIdx.x)/* get block id：2D grid */
#define get_bid() (blockIdx.x+blockIdx.y*gridDim.x)//kernel定义，必须用__global__标识，返回类型必须是void
__global__ void vec_add(double *x, double *y, double *z, int n){int i = get_tid(); //user-defined macro/function//这里的get_tid()是用户定义的一个宏，调用它可以得到当前线程一维的IDif(i<n) z[i] = x[i]+y[i];  //这种情况下启动的线程数大于等于向量的长度，否则会出错
}