Credo ut intelligam

Using NVSHMEM 1、NVSHMEM 简介 针对 CUDA 设备实现的基于 OpenSHMEM 的规范（PGAS 并行编程范式，partitioned global address space） OpenSHMEM 和 NVSHMEM 的区别 NVSHMEM 通信过程中所有的 buffer arguments 必须是对称的 Block 中线程进行 fetch 操作时只能保证数据的 weak ordering（如需强制顺序则要显式调用 nvshmem_fence，其会保证在 fence 之前的所有操作以保证内存一致性） 支持单边通信的方式：通过 Verbs API 或者 UCX 进行 IB 或 RoCE 传输 2、NVSHMEM 优点 支持多 GPU 间直接进行 data movement 或 synchronization，而无需考虑和 CPU 之间的互动之类的操作带来的 overhead 支持通过 composite kernel 实现 comp 和 comm 的 overlap 3、MPI 的问题 可能会引起 shared data 读写的 high locking/atomics overheads Message ordering 会引起 serialization overheads 接收方在发布描述符之前收到数据会引起 protocol overheads 4、NVSHMEM 对应的 addressing model（类似 PGAS） 一个 NVSHMEM job 中的所有 PE 必须同时被初始化，在 job 退出前也需要同时 finalize 通过 CPU-side NVSHMEM allocation API 可以在 GPU 上分配（必须是）大小相同的 shared symmetric memory（而其他方式开辟的内存则属于 private memory） 每个 PE 上的 symmetric memory 对其他 PE 均可见且可操作（通过 NVSHMEM API） 通过 NVSHMEM 分配 API（如 nvshmem_malloc ）返回的对称内存地址对于调用该 API 的 PE 对应的 GPU 来说，是一个可以通过 CUDA API 或者 Load/Store 操作直接访问的有效内存地址。因此如果仅需要对本地 PE 的对称内存进行...

1、NVIDIA最大的敌人不是来自外界，而是它自己：We launched into a wide-ranging discussion of the company’s history. Jensen knows that many of his former employees look back on Nvidia’s beginnings with nostalgia. But he resists overly positive accounts of Nvidia’s start-up period—and his own missteps.  “When we were younger, Tae, we sucked at a lot of things. Nvidia wasn’t a great company on day one. We made it great over thirty-one years. It didn’t come out great,” he said. “You didn’t build NV1 because you were great. You didn’t build NV2 because you were great,” he said, referring to the company’s first two chip designs, both of which were flops that nearly killed Nvidia. “We survived ourselves. We were our own worst enemy.”   There were several more near-death experiences. But each time, amid the stress and the pressure, the company learned from its mistakes. It retained a core of die-hard employees, many of whom remain in the fold to this day. 2、NVIDIA成功的最关键要素——独特的组织架构设计和工作文化：Through these int...

虽说这篇随想的标题叫“用时间对抗时间”，但实际上每个人都明白没有人能赢过时间——从死亡的角度来说，没有人能够让时间停止或倒流，每个人都不可避免得不断逼近自己死亡的时刻。说什么医学或科技能让人永生的那都是扯淡的，这些最多只能延寿。从技术和哲学意义上来看肉体“永生”这件事是很荒谬的，比如瑞典哲学家 Martin Hägglund 认为如果人类获得了永生，反而会发现人生不再有什么是值得为之过活的（ there was nothing left to live for） ，想要战胜死亡的人如果真的美梦成真，那么他们付出的代价比死亡更严重，即失去人生意义。 扯远了。其实我是在一个周末在家撸铁的时候突然想到这件事的。时间是魔法，它让少男少女越发俊美，拥有不断增长的体力精力，却也让中年人的容貌变老、身体机能变得衰弱，活的越来越无精打采；它有时候能够抚平一个人内心的创伤，有时候也能让一个人的痛苦越发剧烈（前段时间看到的一篇文章《 那些父母离世的年轻人 》里提到许多人在多年后仍然无法对年轻时失去父母这件事感到释怀）。那么面对这样的魔法我们是否就无能为力，任由它操控我们的身心状态？这是虚无主义者的借口罢了，我们虽然无法打败时间，但我们可以用时间对抗时间、用魔法对抗魔法。 从锻炼身体的角度来说，五个工作日中，我尽量安排两天晚上能够做下 Zone 2 范围内的有氧（主要就是用壶铃做中等强度的有氧，我的本体感觉告诉我高强度有氧对自己的绝对力量有一定影响，是否真的如此有待后面看看专业人士怎么说），一天晚上做对于神经没有那么大压力的上肢力量训练；同时周末两天分别进行大重量力训（深蹲+硬拉）。作为九零后我已经过了三十岁的年纪了，从正常生理机能的发展来看，心功能、肌肉量和骨密度等指标即将（甚至已经）开始走下坡路，面对这一事实我无法坐以待毙，于是我选择用时间去不断维持甚至提高我的肌肉量、骨骼强度（我发现我最近一段时间体重略微有些涨幅，感觉似乎增肥增肌了）。我知道有一天我的腿还是会变成细狗样，但尽可能延缓身体机能衰退的时间这一行动本身我认为还是挺有意义的，套用里尔克的名言：面对时间的侵袭，有何胜利可言？挺住就是一切。前几天我还收到了 Stronglifts 5X5 训练法创始人 Mehdi Hadim 的一封 newsletter，里面如此回应年轻时不想练腿的人： 除了锻炼，每天保证营养丰富的饮食、拿...

Chap2 Heterogeneous data parallel computing 1、CUDA 代码的结构：The structure of a CUDA C program reflects the coexistence of a host (CPU) and one or more devices (GPUs) in the computer. Each CUDA C source file can have a mixture of host code and device code. By default, any traditional C program is a CUDA program that contains only host code . One can add device code into any source file. 2、CUDA runtime 提供的内存分配和回收接口 cudaMalloc：内存申请函数，用于在设备上开辟一段空间（device global memory）。该函数与cudaMallocManaged 的区别在于后者分配的空间会使用 unified memory 进行自动调度。需要注意该函数的输入的指针类型是 二级指针（void**） ，这样该接口就不会受限于特定数据类型类型的指针 cudaFree：使用 cudaMalloc 开辟完空间后需要使用函数对空间进行释放 float* A_d int size=n * sizeof(float); cudaMalloc((void**)&A_d, size); // A_d为指向device global memory的地址 ... cudaFree(A_d); // 释放分配给A_d的device global memory并放回至available pool cudaMemcpy：用于在主机和设备之间同步数据。第一个入参是目的地址，第二个入参为源地址 3、CUDA Kernel 的运行结构 SPMD 分布式设计模式：single-program multiple-date，指的是 多个计算节点执行相同的程序，但是每个节点处理的数据不同 。SPMD 模型通常用于并行计算，可以将大规模的数据集分成多个小块，由不同的计算节点进行并行处理 blockDi...