原文链接 https://zhuanlan.zhihu.com/p/34587739

一、Gpu的线程结构

要深刻理解kernel，必须要对kernel的线程层次结构有一个清晰的认识。首先GPU上很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。线程两层组织结构如下图所示，这是一个gird和block均为2-dim的线程组织。grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（主要水平方向为x轴），定义的grid和block如下所示，kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。

所以，一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread (1,1)满足：

解释
threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

二、cuda 编码入门

Gpu的编程，kernel快速了解参考如下，对于flash attention 的分块理解有帮助。

CUDA编程入门极简教程

核心：根据计算的特点，构建合理的block 大小（一维，二维，size), 然后计算出grid 的大小。再计算出变量元素x, y和每一个线程x, y对应关系，将每一个计算单元对应上一个线程。

在cpu 中调用kernel 执行函数，然后将结果读取回cpu的内存。

回读结果到CPU 方式：

1. 手动的拷贝

cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);

2. cuda内存托管 + cpu, gpu的同步 

// 同步device 保证结果能正确访问cudaDeviceSynchronize();

// 同步device 保证结果能正确访问 cudaDeviceSynchronize();

CUDA编程入门极简教程