CUDA 编程模型
线程→块→网格
网格如何划分为块
dim3 dimGrid(x,y,z);定义了一个z纵y行x列的网格，即x,y,z是从“变化最快的”维度开始的，对比等价C语言写法应该是mat[z][y][x]
想想C语言中复合类型是怎么阅读的-由内而外。
//大小为10的数组，它的每个元素都是大小为20的数组
//这些数组的元素是含有30个整数的数组
int arr[10][20][30] = {0}；//将所有元素初始化为0因此，对于二维情况
row = blockIdx.y * blockDim.y + threadIdx.y // vertical cord
col = blockIdx.x * blockDim.x + threadIdx.x // horizontal cord
mat[row][col] = xxx; //这样，同一个warp内的连续（变化最快，用x来区分）线程访问了连续的内存
行主序 / 列主序
row-major order
描述了多维数组在线性存储中的存储方式
Row- and column-major order - Wikipedia
arch
GPU中有很多SM，和一个大DRAM。
每个SM中有一些核。
多个块可以被调度到同一个SM
块可以被划分为线程束，线程束一般是32个线程，线程束是SM内部的调度单位。
SM中的核被分组到不同处理块，线程束在处理块上被执行
35
C++ Rules of Three, Five & Zero | Sonar
c++ - Is there any case where a return of a RValue Reference (&&) is useful? - Stack Overflow
多线程环境下waitpid的逻辑并没有变
总有一个线程会先执行_exit
posix_spawn在最新glibc中由clone和exec实现
条件变量比信号量更加通用，更加好写
条件变量在上锁状态检查全局状态，有全局信息不容易错
而信号量只观察少数局部状态来试图完成正确的同步
如哲学家吃饭问题中，信号量只关心周围两个叉子，而条件变量使用时有全局叉子信息
各种并发bug
数据竞争-行为不确定
死锁
原子性违反和顺序违反
CPU选择了性能（延迟？）而不是能效（吞吐量？）
寻找指令级并行机会，需要额外的电路 能耗 热量
会撞上功耗墙
需要换思路了 SIMD和多/异构处理器