GPS: A Global Publish Subscribe Model for Multi-GPU Memory Management

<ul>
 <li><a href="https://www.microarch.org/micro54/">MICRO’21</a> Best Paper Nominee</li>
 <li><a href="https://dl.acm.org/doi/10.1145/3466752.3480088">PDF</a></li>
</ul>

<h2 id="background-and-motivation">BACKGROUND AND MOTIVATION</h2>

<ul>
 <li>GPU：高带宽、适合并行的硬件架构
 <ul>
 <li>SIMD + DSA（Tensor Core）</li>
 </ul>
 </li>
 <li>多 GPU 可提供更高性能，但充分管理硬件资源难度大</li>
 <li>主要困难之一：本地和远程存储带宽相差数量级，互连带宽限制可扩展性
 <ul>
 <li>4 卡，无限带宽（理想情况）、PCIe 6.0（预测）、PCIe 3.0 相对于单卡的性能分别为：
 <ul>
 <li>$3\times, 2\times, 0.7\times$</li>
 </ul>
 </li>
 </ul>
 </li>
 <li>当前多 GPU 系统中有效管理数据的技术均存在不足</li>
</ul>

<ul>
 <li>主机端启动 GPU 之间的 memcpy
 <ul>
 <li>需要精细调整以实现计算通信的流水</li>
 <li>MPI、NCCL 等通信框架可直接操作 GPU 缓冲区
 <ul>
 <li>计算和通信的重叠仍然需要手动调整</li>
 <li>学习成本增加，进一步加重程序员的负担</li>
 </ul>
 </li>
 </ul>
 </li>
 <li>点对点对等传输
 <ul>
 <li>GPU 线程可以直接访问其他 GPU 的物理内存</li>
 <li>细粒度执行，计算通信重叠，启动开销低（只需要一方即可通信，不用同步）
 <ul>
 <li><a href="https://developer.nvidia.com/nvshmem">NVSHMEM</a> 较之 MPI 更高效</li>
 </ul>
 </li>
 <li>严重受限于互连带宽，读远程内存时延迟高</li>
 </ul>
 </li>
</ul>

<ul>
 <li>通过统一寻址内存（Unified Memory, UM）
 <ul>
 <li>UM 提供系统中任何处理器均可访问的单一且统一虚拟内存地址空间</li>
 <li>UM 以隐式、和程序员无关的方式支持 GPU 之间的数据移动
 <ul>
 <li>使用页面错误和迁移机制在 GPU 之间执行数据传输</li>
 <li>页面错误的性能开销巨大，只是看起来美好</li>
 </ul>
 </li>
 <li>UM 可以通过显式提示的方式减少页面错误
 <ul>
 <li>仍然实现细粒度数据共享的流水线预取和计算</li>
 </ul>
 </li>
 <li>不支持至少有一个 writer 和多个 reader 的页面复制
 <ul>
 <li>对有多个 reader 的页面的写将页面“折叠”到单个 GPU（通常是 writer）会触发昂贵的 TLB</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<ul>
 <li>避免远程访问瓶颈的一种方法：
 <ul>
 <li>在生成数据后，提前将数据从生产 GPU 传输到消费 GPU</li>
 <li>在需要时直接从本地内存读取数据</li>
 </ul>
 </li>
 <li>主动传输有助于实现强大的可伸缩性
 <ul>
 <li>提供了更多的机会使数据传输与计算重叠</li>
 <li>改进了局部性，并确保关键路径负载享有更高的本地内存带宽</li>
 </ul>
 </li>
 <li>只有一部分 GPU 将消耗数据的情况下，执行广泛的全对全传输会浪费 GPU 间的带宽</li>
</ul>

<ul>
 <li>GPS：一个全局发布-订阅的多显卡内存管理模型，提供了一组架构增强功能
 <ul>
 <li>订阅管理完全由发布-订阅处理单元处理</li>
 <li>自动跟踪哪些 GPU 订阅共享内存页</li>
 <li>主动向所有订户广播存储，使订户能够以高带宽从本地内存读取数据</li>
 <li>通过驱动程序，在每个订阅 GPU 上本地复制这些页</li>
 <li>发布-订阅可以解耦</li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 的读操作总是从本地返回，而 GPS 的写操作则广播给所有用户
 <ul>
 <li>传统的加载-存储根据物理内存位置导致本地或远程访问，带来延迟</li>
 <li>在写路径而不是读路径上执行远程访问会隐藏延迟
 <ul>
 <li>远程存储不会暂停执行</li>
 </ul>
 </li>
 <li>可以进行进一步优化
 <ul>
 <li>对写操作调度、组合，有效地使用互连带宽</li>
 <li>不违反内存模型</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<ul>
 <li>一个简单的发布-订阅模型示例
 <ul>
 <li>生成数据的发布者</li>
 <li>请求特定更新的订阅者</li>
 <li>发布-订阅处理单元</li>
 </ul>
 </li>
</ul>

<ul>
 <li>依赖主动式远程存储的发布-订阅模型的主要挑战
 <ul>
 <li>应由哪个 GPU 接收数据？</li>
 <li>应该何时传输存储？</li>
 </ul>
 </li>
</ul>

<h2 id="gps-architectural-principles">GPS ARCHITECTURAL PRINCIPLES</h2>

<ul>
 <li>GPS 的设计目标
 <ul>
 <li>不破坏与 GPU 编程模型、内存一致性模型的向后兼容性
 <ul>
 <li>提供一致的性能增益</li>
 </ul>
 </li>
 <li>GPS 可以轻松被集成到应用程序中
 <ul>
 <li>只需极少的代码或概念更改</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 背后的关键创新
 <ul>
 <li>一系列架构增强，维护订阅信息</li>
 <li>旨在协调主动式 GPU 间通信以实现高性能</li>
 <li>仅向需要的 GPU 传输数据</li>
 <li>在使用过程中实现本地访问</li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 地址空间
 <ul>
 <li>扩展传统多 GPU 共享虚拟地址空间
 <ul>
 <li>使 GPS 可以通过简单、直观的 API 集成到其应用程序中</li>
 <li>只需对分配和订阅管理进行微小的更改，不需要修改为 UM 编写的 GPU 内核</li>
 <li>很容易集成到现有的多 GPU 编程框架上</li>
 </ul>
 </li>
 <li>在 GPS 地址空间中的分配在所有订阅 GPU 的物理内存中都有本地副本</li>
 <li>语法相同，但底层行为不同
 <ul>
 <li>拦截每个读写操作并将副本转发给每个订户的本地副本</li>
 <li>读操作可以在完全本地带宽下执行，无需消耗互连带宽</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 提供手动和自动机制来管理页面订阅
 <ul>
 <li>手动：允许用户通过订阅/取消订阅 API 显式指定订阅信息</li>
 <li>自动：可以通过以下两种方式之一执行：
 <ul>
 <li>默认情况下订阅，即不分青红皂白地全部订阅，然后是取消订阅阶段</li>
 <li>默认情况下不订阅，即 GPU 仅在向某个页面发出第一个读取请求时才订阅该页面</li>
 <li>一旦捕获到共享者信息，共享者信息就会反馈到用于协调 GPU 间通信的订阅跟踪机制中</li>
 </ul>
 </li>
 <li>订阅是两种机制对 GPS 的提示，而不是正确执行应用程序的功能要求。
 <ul>
 <li>如果 GPU 向它不是订户的页面发出加载，它不会出错</li>
 <li>对远程 GPU 执行点对点加载会造成性能损失，但不会破坏功能正确性</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<ul>
 <li>合并写操作
 <ul>
 <li>GPS 在将存储转发到其他 GPU 之前，会将存储主动合并到 GPS 页面
 <ul>
 <li>节省了宝贵的 GPU 间带宽</li>
 </ul>
 </li>
 <li>对于作用域访问，像传统 GPU 那样简单地处理它们
 <ul>
 <li>GPU 内存模型只要求在执行作用域同步（如内存围栏）时，其他 GPU 可以看到写操作</li>
 <li>所有作用域访问都被发送到单个一致性点并在那里执行</li>
 <li>仅在多 GPU 上并发启动的 grid 需要通过内存显式同步时使用
 <ul>
 <li>通常这类访问很少，所以无问题！</li>
 </ul>
 </li>
 </ul>
 </li>
 <li>对于非作用域访问仅维护必需的读写排序行为
 <ul>
 <li>同一 GPU 到同一地址的读写以相同的顺序到达</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<h2 id="gps-programming-interface">GPS PROGRAMMING INTERFACE</h2>

<ul>
 <li><code class="language-plaintext highlighter-rouge">cudaMallocGPS</code>：内存分配
 <ul>
 <li><code class="language-plaintext highlighter-rouge">cudaMalloc</code>的代替</li>
 <li>在 GPS 虚拟地址空间中分配内存，并在至少一个 GPU 中使用物理备份。</li>
 <li>在分配时传递一个可选参数，以指示将为该区域显式管理订阅</li>
 <li>否则，GPS 将执行自动订阅管理</li>
 <li>使用<code class="language-plaintext highlighter-rouge">cudaFree</code>释放</li>
 </ul>
 </li>
</ul>

<ul>
 <li><code class="language-plaintext highlighter-rouge">cuMemAdvise</code>：内存订阅提示
 <ul>
 <li><code class="language-plaintext highlighter-rouge">CU_MEM_ADVISE_GPS_SUBSCRIBE</code>：手动订阅</li>
 <li><code class="language-plaintext highlighter-rouge">CU_MEM_ADVISE_GPS_UNSUBSCRIBE</code>：取消订阅</li>
 </ul>
 </li>
 <li>自动分析并订阅
 <ul>
 <li><code class="language-plaintext highlighter-rouge">cuGPSTrackingStart</code>：开始观察</li>
 <li><code class="language-plaintext highlighter-rouge">cuGPSTrackingStop</code>：结束观察并分析订阅</li>
 <li>分析不需要精确以保持正确性
 <ul>
 <li>GPU 仍可以访问未订阅的页面</li>
 <li>但这样的访问将以较低的性能执行</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<div class="language-cpp highlighter-rouge"><div class="highlight"><pre class="highlight"><code>__global__ void mvmul(float *mat, float *invec, float *outvec /* ... */) {
 /* ... */
 for (int i = 0; i < mat_dim; ++i)
 sum += mat[tid * mat_dim + i] * invec[i];
 outvec[tid] = sum;
 /* ... */
}
</code></pre></div></div>

<div class="language-cpp highlighter-rouge"><div class="highlight"><pre class="highlight"><code>cudaMallocGPS(&mat, mat_dim *mat_dim_size);
cudaMallocGPS(&vec1, mat_dim_size);
cudaMallocGPS(&vec2, mat_dim_size);
cudaMemset(vec2, 0, mat_dim_size);
for (int iter = 0; iter < MAX_ITER; ++iter) {
 if (iter == 0)
 cuGPSTrackingStart();
 for (int device = 0; device < num_devices; ++device) {
 cudaSetDevice(device);
 mvmul<<<num_blocks, num_threads, stream[device]>>>(mat, vec1,
 vec2 /* ... */);
 mvmul<<<num_blocks, num_threads, stream[device]>>>(mat, vec2,
 vec1 /* ... */);
 }
 if (iter == 0)
 cuGPSTrackingStop();
}
</code></pre></div></div>

<h2 id="architectural-support-for-gps">ARCHITECTURAL SUPPORT FOR GPS</h2>

<ul>
 <li>一种可能的 GPS 硬件实现
 <ul>
 <li>GPU 页表条目（PTE）中增加一位，以指示虚拟内存页是否是 GPS 页</li>
 <li>一个新的硬件单元，负责将写操作传播到订阅了特定页面的 GPU</li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 必须支持以下基本内存操作
 <ul>
 <li>传统的加载、存储和原子操作</li>
 <li>GPS 加载</li>
 <li>GPS 存储和原子操作</li>
 </ul>
 </li>
</ul>

<ul>
 <li>GPS 硬件单元和扩展
 <ul>
 <li>页表支持</li>
 <li>合并远程写</li>
 <li>GPS 地址转换</li>
 <li>访问跟踪单元
 <ul>
 <li>用于 Profile</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<h2 id="experimental-methodology">EXPERIMENTAL METHODOLOGY</h2>

<ul>
 <li>对 NVAS（NVIDIA Architectural Simulator）进行修改
 <ul>
 <li>NVAS 在 SASS 级别进行驱动</li>
 </ul>
 </li>
 <li>在实际硬件上使用工具 NVBit，跟踪包含
 <ul>
 <li>CUDAAPI 事件</li>
 <li>GPU 内核指令</li>
 <li>访问的内存地址</li>
 </ul>
 </li>
</ul>

<table>
 <thead>
 <tr>
 <th style="text-align: center">Application</th>
 <th style="text-align: center">Description</th>
 <th style="text-align: center">Predominant Communication Pattern</th>
 </tr>
 </thead>
 <tbody>
 <tr>
 <td style="text-align: center">Jacobi</td>
 <td style="text-align: center">Iterative algorithm that solves a diagonally dominant system of linear equations</td>
 <td style="text-align: center">Peer-to-peer</td>
 </tr>
 <tr>
 <td style="text-align: center">Pagerank</td>
 <td style="text-align: center">Algorithm used by Google Search to rank web pages in their search engine results</td>
 <td style="text-align: center">Peer-to-Peer</td>
 </tr>
 <tr>
 <td style="text-align: center">SSSP</td>
 <td style="text-align: center">Shortest path computation between every pair of vertices in a graph</td>
 <td style="text-align: center">Many-to-many</td>
 </tr>
 <tr>
 <td style="text-align: center">ALS</td>
 <td style="text-align: center">Matrix factorization algorithm</td>
 <td style="text-align: center">All-to-all</td>
 </tr>
 </tbody>
</table>

<table>
 <thead>
 <tr>
 <th style="text-align: center">Application</th>
 <th style="text-align: center">Description</th>
 <th style="text-align: center">Predominant Communication Pattern</th>
 </tr>
 </thead>
 <tbody>
 <tr>
 <td style="text-align: center">CT</td>
 <td style="text-align: center">Model Based Iterative Reconstruction algorithm used in CT imaging</td>
 <td style="text-align: center">All-to-all</td>
 </tr>
 <tr>
 <td style="text-align: center">B2rEqwp</td>
 <td style="text-align: center">3D earthquake wave-propogation model simulation using 4-order finite difference method</td>
 <td style="text-align: center">Peer-to-peer</td>
 </tr>
 <tr>
 <td style="text-align: center">Diffusion</td>
 <td style="text-align: center">A multi-GPU implementation of 3D Heat Equation and inviscid Burgers’ Equation</td>
 <td style="text-align: center">Peer-to-peer</td>
 </tr>
 <tr>
 <td style="text-align: center">Hit</td>
 <td style="text-align: center">Simulating Homogeneous Isotropic Turbulence by solving Navier-Stokes equations in 3D</td>
 <td style="text-align: center">Peer-to-peer</td>
 </tr>
 </tbody>
</table>

<h2 id="experimental-results">EXPERIMENTAL RESULTS</h2>

<ul>
 <li>Limitations of the GPS approach
 <ul>
 <li>订阅以页（64KB）为单位
 <ul>
 <li>即使只有一位更新，也需要传输整页</li>
 </ul>
 </li>
 <li>GPU 通常以缓存块为粒度进行互联
 <ul>
 <li>导致传输不必要的字节</li>
 </ul>
 </li>
 <li>写合并技术可能不会合并时间上间隔较远的操作
 <ul>
 <li>导致效率降低</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>

<h2 id="conclusion">CONCLUSION</h2>

<ul>
 <li>GPS：一种全新的软硬件共同设计的多 GPU 内存管理系统扩展
 <ul>
 <li>通过发布-订阅模式提高多 GPU 系统性能</li>
 </ul>
 </li>
 <li>GPS 有简单直观的编程模型扩展，可以自然地被集成到应用程序中</li>
 <li>GPS 提供了显著的性能改进
 <ul>
 <li>在 4 个 NVIDIA V100 GPU 和几种互连架构的模型上进行评估
 <ul>
 <li>GPS 在 1 个 GPU 上的平均加速比为 3.0</li>
 <li>性能比当前最佳可用多 GPU 编程模式好 2.3 倍。</li>
 </ul>
 </li>
 <li>在类似的 16 GPU 系统上
 <ul>
 <li>GPS 捕获了 80%的可用性能</li>
 <li>大大领先于其他多 GPU 编程模型。</li>
 </ul>
 </li>
 </ul>
 </li>
</ul>