Credo ut intelligam

 一、双缓冲模式 双缓冲模式是一种常见的提升数据发送效率的方法，数据的发送和接收往往有延迟。如果只用一块缓冲区，必须等通信完成后才能复用这块内存。而使用双缓冲则让通信任务和数据准备任务同时进行，有效隐藏通信延迟，提高整体吞吐。该方法涉及两块缓存，第一块缓存用于当前数据收发时，也会对第二块缓存会进行清理以准备下一轮的数据发送。例如下面的 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 方法，该方法会根据传入的原始数据指针创建一个引用（而非复制） 指向...

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...

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...