大模型与AI底层技术揭秘（37）绞刑架下的报告

在上期，我们留下了一个问题：在warp中，多个CUDA核需要同步的时候需要怎么做呢？

我们先来看一个故事。

1939年，纳粹德国的铁蹄碾碎了捷克美好的山河。捷克共产党人伏契克与他的战友们被迫转入地下斗争。1942年，由于叛徒的出卖，伏契克和战友们被盖世太保逮捕了。

在纳粹的监狱中，伏契克受到了非人的折磨，也不知道何时自己的家乡能重获解放。但是，伏契克依然相信，纳粹终将失败！他在狱中建立了临时党组织，并设法让狱警中的一名捷克人帮助他获取了纸和笔，写下了名著《绞刑架下的报告》。1945年，苏联红军攻克柏林，将伏契克的妻子救出监狱，但伏契克没有看到这天。他在受尽了种种肉体和精神折磨后，在1943年9月8日，化作了镰刀锤子旗下的一颗金星。

在茫茫的黑暗中，伏契克燃烧了自己的生命，为狱中的战友们照亮了通往胜利的曲径。在《绞刑架下的报告》中，伏契克描述了这个坚强勇敢，肝胆相照的战斗集体。反法西斯战争胜利后，《绞刑架下的报告》也成为了不朽的经典。

与此同时，在隔着小半个地球的远东战场，却出现了反法西斯战争中非常无耻的一幕。一名军长，在二战胜利的前夜，由于贪生怕死，在部队仍有战斗力，后方一直进行空投补给的情况下，率部投降鬼子，最终投降的官兵大部分被杀害，这名军长却苟且偷生到1983年才病死。

为什么在反法西斯战争中，二者的表现有天壤之别呢？

重要的一个原因就是，共产党人是严格的组织的，大家严格保持一致，而这个反动军长所在的党派，本质上是封建组织，其价值观是扭曲的。

这种与全人类反法西斯步调背离的卑鄙小人，最终还是难以避免被钉在全人类的历史耻辱柱上，而一切为这种人洗白的所谓作家或导演，也将被扫入历史的垃圾堆。

小H对比了两个故事，也理解了，在GPU这样高度并行的SIMT处理器中，各个CUDA核心同步的重要性。

CUDA框架提供了这一核心同步机制，也就是函数__syncthreads()。

当一个thread(也就是GPU中的CUDA核心）执行到__syncthreadas()时，这个thread会看它所在的block（也就是GPU中的线程束warp）内的其他所有threads情况，如果发现还有其他threads没有执行到这个位置，则这个thread等待其他threads。直到block中所有active的threads都执行到此，接着向下执行。

注意到线程同步有可能产生死锁，如：

if (func()){
    __syncthreads();
}
else{
    __syncthreads();
}

这样会导致不同分支的threads互相等待，谁也等不到谁。

因此，只有整个线程块的分支全部合流的时候，才可以使用__syncthreads()，否则一定会造成死锁。

除了__syncthreads以外，还有另外三个在CPU侧执行的同步函数：

cudaDeviceSynchronize：它将阻塞CPU端线程的执行，直到每个CPU端线程都执行完毕了GPU侧的核函数。

cudaThreadSynchronize：它和cudaDeviceSynchronize基本相同，但最新版本的CUDA已经不建议使用。

cudaStreamSynchronize：它是cudaDeviceSynchonize的增强。主要的增强就是可以传入cudaStreamID参数，只阻塞指定流ID的线程。

有CPU下多线程/多进程开发经验的同学有可能会回忆起一个概念：内存屏障（memory barrier，也称为内存栅栏）

所谓的内存屏障，指的是：

内存屏障之前的所有写操作都要写入内存；内存屏障之后的读操作都可以获得同步屏障之前的写操作的结果。

类似地，GPU中也有内存屏障机制。

在CUDA库中，提供了__threadfence()等机制实现内存屏障。在调用__threadfence后，可以保证该线程在调用前，对全局存储器或共享存储器的访问已经全部完成，执行结果对线程网格中的所有线程是可见的。易言之，当一个线程运行__threadfence后，线程在这个时刻之前，对于存储器的读取或写入，对所有网格（warp）内的线程都是有效的。

我们在理解了GPU的线程调度和线程同步机制以后，就可以深入理解GPU的存储器组织和指令了。

请看下期。