博文

硬芯与硬心

前段时间在盒马上买了几袋零脂玉米面，其配料表很干净，只有玉米粉、饮用水和食用盐，每百克钠含量是 7%，在面条里算比较低的了。拿到手后看了看食用方法，发现有 2 种：第一种是提前把面条在凉水里浸泡 2 小时以上，然后再下锅煮；第二种是直接冷水下锅，开锅后再煮 2-3 分钟然后关火继续焖几分钟。虽然两种做法略有不同，但它们有个共同点：待面条无硬芯即可食用。虽然看了说明，但第一次煮的时候还是大意了，以为用筷子稍用力夹就断就可以捞出来吃了，没想到吃起来却很“劲道”——当然这只是未将其再煮软一些的借口罢了——嚼起来的时候仍能感受到来自面条内“芯”的抵抗。于是后来几次就学乖了，煮的时间甚至比食用方法里建议的时间还更久一些，此时煮出来的面条口感正好符合自己的要求，看来“待面条无硬芯即可食用”这句话得划重点。 硬芯是这款玉米面的特点，而非它的错误，但人却不一样。如果当一个人内心刚硬的程度是玉米面的好多倍，面对他人正确且善意的指责不去虚心接受并改正，以为依然可以我行我素，那这样的硬心往往会带来悲剧。埃及记中的法老，上帝透过摩西不断降灾，为了消灾法老表面上次次答应摩西的条件却又次次“心里刚硬”而毁约，最终落得个埃及地所有人和牲畜生的长子必定死的下场，而这些人和动物本可以不死；西底家登基做犹大王的时候“行耶和华他神眼中看为恶的事。先知耶利米以耶和华的话劝他，他仍不在耶利米面前自卑...强项硬心，不归服耶和华以色列的神 ”（代下 36:12-13），将耶路撒冷变成了污秽之地，注意经文中说“所以”迦勒底人来攻击他们，焚烧圣殿，毁坏耶路撒冷，最后的下场是“地土荒凉”（36:21）；当施洗约翰传悔改的道时，法利赛人和撒都该人也想一起凑热闹，但约翰却说他们是“毒蛇的种类”，因为这些人自诩亚伯拉罕的后裔——“有亚伯拉罕在我们怕什么！”这也是硬心而非悔改的一种表现，如果继续选择硬心，那么结局就是被看下来丢在火里（太 3:10）。 如果你留意世界上每天发生的许多悲剧，会发现上帝的审判并非只会跟着世界末日一起到来，因为一个人人做自己任意想要做的世界就是上帝给人的审判，正如路易斯在《梦幻巴士》（*The Great Divorce*）说：“归根究底只有两种人。一种人对上帝说：‘愿你的旨意成就’；另一种人上帝最后对他们说：‘愿你的旨意成就’”。硬心的人认为“人的所愿就是他的天堂”，殊不知“人的所愿就是他的地狱...
发表评论
阅读全文

DeepEP原理分析（一）：low latency模式特点总结

图片
 一、双缓冲模式 双缓冲模式是一种常见的提升数据发送效率的方法，数据的发送和接收往往有延迟。如果只用一块缓冲区，必须等通信完成后才能复用这块内存。而使用双缓冲则让通信任务和数据准备任务同时进行，有效隐藏通信延迟，提高整体吞吐。该方法涉及两块缓存，第一块缓存用于当前数据收发时，也会对第二块缓存会进行清理以准备下一轮的数据发送。例如下面的 low latency dispatch 代码中第一块缓存为 buffer，第二块为 next_buffer，它们均为 LowLatencyBuffer 结构体，位于显存（rdma_buffer_ptr ）上 当 dispatch kernel 利用多个 SM 和 buffer 进行数据收发时，最后一个 warp 中的第一个 SM 也会对 next_buffer 进行清零操作，这样当下一次进行 combine 时可以直接复用刚刚完成的 dispatch 过程中准备好的 next_buffer，而 combine 过程同样也会为下一次的 dispatch 准备好 buffer。buffer 和 next_buffer 见的切换通过“异或”操作进行（^= 1），每进行一次异或操作 low_latency_buffer_idx 的值就会从 0 变为 1 或 1 变成 0，这样每轮 dispatch 和 combine 的过程就能错开利用两块 buffer 二、zero copy 原理 在 test_low_latency 测试脚本中有个叫做 test_func 的函数，该函数允许调用者手动设置是否使用 zero copy 模式来进行 combine。顾名思义 zero copy 指在 combine——也就是各专家将 dispatch 过程中收到的 token 计算完后的结果发回 token 所在的原 rank——该过程中无需将 token 计算结果拷贝至 RDMA 发送缓冲区，而是直接将让 combine 发送缓冲区（combine_rdma_send_buffer_data_start）指向计算后结果（simulated_gemm_x） 其关键作用的便是 get_next_low_latency_combine_buffer 函数，里面会使用 from_blob 方法，该方法会根据传入的原始数据指针创建一个引用（而非复制） 指向...
发表评论
阅读全文

NVSHMEM官方文档部分内容总结

图片
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 的对称内存进行...
发表评论
阅读全文

The Nvidia Way书摘

图片
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，里面如此回应年轻时不想练腿的人： 除了锻炼，每天保证营养丰富的饮食、拿...
发表评论
阅读全文

Programming Massively Parallel Processors部分章节重点摘录

图片
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...
发表评论
阅读全文