第6章精讲：MUSA 编程基础（编译链接 / 索引 / 错误处理 / 归约实战）

一、设备管理：一切资源分配的前提

设备是 MUSA 异构计算的硬件核心，所有计算与存储资源均依托设备实现——设备管理是后续资源分配、代码执行（如核函数启动）的前提。MUSA 提供 3 种方法来查询和管理 MT GPU 设备：MUSA-runtime API、MUSA 环境变量、mthreads-gmi 工具。

1.1 MUSA-runtime API：查询和设置 GPU

• musaError_t musaGetDeviceCount(int *count)：获取当前环境的可用 GPU 数量。
• musaError_t musaSetDevice(int device)：在多 GPU 环境中选择使用的 MUSA 设备（默认使用 GPU0，编号范围为 0～count-1）。后续创建的上下文、流、内存等 GPU 资源都会在设置的设备上。
• musaError_t musaGetDeviceProperties_v2(struct musaDeviceProp *prop, int device)：查询当前环境中每一个设备的属性。musaDeviceProp 结构体包含设备核心参数，详见原书第 6 章代码文件 musaDevicePropDefinition.cpp；原书第 6 章代码文件 getDeviceBasicProperties.cpp 展示了如何获取所有设备的基本属性。

通过 musaInfo 命令，用户可全面获取摩尔线程 GPU 设备的硬件配置与运行时状态信息，以便进行性能优化与资源规划。编译命令为：

g++ musaInfo.cpp -o musaInfo -I/usr/local/musa/include/ -L /usr/local/musa/lib -lmusart

1.2 选择最优 GPU：以 multiProcessorCount 为依据

当存在多个 GPU 设备时，可以通过 multiProcessorCount（表示处理器数量）来选择算力最优的 GPU。完整示例代码如下：

#include <musa_runtime.h>
#include <cstdio>

int selectBestGPU()
{
  int numDevices       = 0;
  int maxMultiprocessors = 0;
  int bestDevice        = 0;
  musaError_t err = musaGetDeviceCount(&numDevices);
  if (err != musaSuccess || numDevices == 0)
  {
      fprintf(stderr, "No musa devices found or failed to get device count.
");
      bestDevice = -1;
  }
  if (bestDevice != -1)
  {
      for (int device = 0; device < numDevices; ++device)
      {
          musaDeviceProp props;
          if (musaGetDeviceProperties(&props, device) != musaSuccess)
          {
              fprintf(stderr, "Failed to get properties for device %d
", device);
              continue;
          }
          if (props.multiProcessorCount > maxMultiprocessors)
          {
              maxMultiprocessors = props.multiProcessorCount;
              bestDevice         = device;
          }
      }
  }
  return bestDevice;
}

int main(int argc, char** argv)
{
    int devId = selectBestGPU();
    fprintf(stdout, "The best performce gpu is GPU%d
", devId);
    return 0;
}

1.3 MUSA 环境变量：不改代码选设备

MUSA Runtime 提供了环境变量，用于在不更改代码和不重新编译的情况下，选择 GPU 设备。

• MUSA_DEVICE_ORDER：控制 MUSA 设备枚举顺序是按照性能优先（FASTEST_FIRST）还是按照 PCIe 总线 ID（PCI_BUS_ID）优先。默认按照性能优先方式。
• MUSA_VISIBLE_DEVICES：控制 MUSA 可见的设备列表，例如：

  # 只有 device 0 可见，musaGetDeviceCount 获取的 device count 为 1
  export MUSA_VISIBLE_DEVICES=0
  # 只可见前 4 张卡
  export MUSA_VISIBLE_DEVICES=0,1,2,3
  # 将设备号从后往前枚举
  export MUSA_VISIBLE_DEVICES=7,6,5,4,3,2,1,0

1.4 mthreads-gmi 工具：命令行快速查看

直接执行 mthreads-gmi 命令可以获取 GPU 名称、驱动版本、使用率等信息：

$ mthreads-gmi
Tue Jan 20 06:04:46 2026
------------------------------------------------------------
mthreads-gmi:2.2.0 Driver Version:3.3.3-server
------------------------------------------------------------
ID Name |PCIe |%GPU Mem
Device Type |Pcie Lane Width |Temp MPC Capable | ECC Mode
--------------------------------------------------------------+
0 MTT S5000 |00000000:01:00.0 |0% 0MiB(65536MiB)
Physical |16x(16x) |55C YES | N/A
--------------------------------------------------------------
--------------------------------------------------------------
Processes:
ID PID Process name GPU Memory Usage
--------------------------------------------------------------+
No running processes found
--------------------------------------------------------------

还有一些 GPU 属性可以通过 mthreads-gmi -q 命令查询到，如序列号、BIOS 版本等：

$ mthreads-gmi -q
Timestamp : Tue Jan 20 06:07:50 2026
Driver Version : 3.3.3-server
Attached GPUs : 1
GPU0 00000000:01:00.0
Product Name : MTT S5000
Product Brand : MTT
GPU UUID : 75b26540-1173-b8d0-9464-52121455d780
Serial Number : E47S02MY13VG0003
DeviceType : Physical
MPC Capable : YES
MTBios Version : 4.3.28
...
Aggregate
DRAM Correctable : N/A
DRAM Uncorrectable : N/A
Retired Pages
Single Bit ECC : N/A
Double Bit ECC : N/A
Pending Page Blacklist : N/A
Performance State : N/A

二、模块管理：内核代码的加载与卸载

模块（Module）是代码与数据一同加载的数据集合，类似于 Windows 操作系统中的 DLL 和 Linux 操作系统中的 DSO（动态共享对象）。MUSA 中没有和对象文件一样可以被整合到模块中的中间结构，取而代之的是 mcc 直接生成的可以被加载为 MUSA Module 的文件，即面向特定 MP 版本的 fatbin 文件。

一旦 MUSA Module 被加载，应用程序便可以查询包含在其中的资源：全局变量、函数（内核）、纹理引用（Texture Reference）。关键查询函数：

• muModuleGetGlobal()：返回模块中的指针与符号大小。
• muModuleGetTexRef()：返回模块中声明的纹理引用大小。
• muModuleGetFunction()：返回模块中声明的内核函数。

所有的资源在模块加载时就被创建，所以查询函数不会因缺少资源而查找失败。

2.1 模块编译：fatbin 文件

模块通过调用 mcc 创建。fatbin 文件可以被编译至一个或多个 GPU 架构，且它们不需要即时编译，因此可以非常快地加载。若不指定特定的 GPU 架构，则默认编译为包含全部 GPU 架构的 fatbin 文件。

# 默认编译为兼容所有 GPU 架构的 fatbin 文件
mcc -mtgpu -c --cuda-device-only kernel.mu -o kernel.fatbin

# 编译至特定 GPU 架构（如 mp_31）的 fatbin 文件
mcc -mtgpu -c --cuda-device-only --offload-arch=mp_31 kernel.mu -o kernel.fatbin

环境变量 MUSA_MODULE_LOADING 用于指定应用程序模块的加载模式，包括 DEFAULT（默认）、LAZY（延迟）、EAGER（立即）：

• 当设置为 EAGER 时，所有内核和数据会在调用 muModuleLoad 或 muModuleLoadData API 时从 fatbin 文件中完全加载。
• 当设置为 LAZY 时，特定内核的加载会延迟到调用 muModuleGetFunction 或 muKernelGetFunction API 获取 MUfunc 句柄时；而 fatbin 文件中的数据会在加载该 fatbin 文件中的第一个内核或首次访问变量时加载。

2.2 MUSA Runtime API 模块管理

像 MUSA 上下文一样，MUSA Runtime 不显式支持模块，模块只在 MUSA Driver API 中可用。所有的模块在 MUSART 初始化时加载。对于拥有大量 GPU 代码的应用程序，选择 MUSA Driver API 而不选择 MUSA Runtime API 的一个首要原因是：前者可以显式地管理模块的加载和卸载。

__global__ void kernel() {
    // 内核代码
}
int main() {
    kernel<<<1,1>>>();
    musaDeviceSynchronize();
    return 0;
}

2.3 Driver API 模块管理：精细粒度控制

应用程序使用 Driver API 进行模块管理的优势是可以对模块的粒度进行控制。例如，它可以作为资源嵌入可执行文件中，在运行时加密或生成，但是使用 fatbin 文件的过程必须显式实现。

__global__ void kernel()
{
  // 内核代码
}
// main 函数的示例代码
#include <musa.h>
int main()
{
  MUmodule module;
  MUfunction kernel;

  // 初始化 Driver API
  muInit(0);

  // 创建上下文（绑定设备 0）
  MUdevice device;
  muDeviceGet(&device, 0);
  MUcontext context;
  muCtxCreate(&context, 0, device);

  // 加载模块
  muModuleLoad(&module, "kernel.fatbin");

  // 获取内核函数
  muModuleGetFunction(&kernel, module, "kernel");

  // 启动内核
  void* args[] = { /* 参数 */ };
  muLaunchKernel(kernel,
    1, 1, 1,   // 网格维度
    1, 1, 1,   // 线程块维度
    0, NULL,   // 共享内存大小和流句柄
    args, NULL);
  muCtxSynchronize();

  // 资源清理（卸载模块并销毁上下文）
  muModuleUnload(module);
  muCtxDestroy(context);
  return 0;
}

2.4 延迟加载：降低初始化开销

延迟加载是指将 MUSA 模块和内核的加载时间从程序初始化阶段推迟到接近内核实际执行时。如果程序未使用其包含的所有内核，则部分内核会被不必要地加载——这种情况在使用库时非常常见，大多数程序通常只会调用库中的少量内核。通过延迟加载，程序可以仅加载实际会用到的内核，从而减少初始化时间，并降低 GPU 内存和主机内存的开销。启用延迟加载需将环境变量 MUSA_MODULE_LOADING 设置为 LAZY。

具体优化分为两部分：

1. MUSA Runtime 的模块加载优化：程序初始化时不再加载所有模块（包含托管变量的模块除外），每个模块会在首次使用其变量或内核时加载。此优化仅影响 MUSA Runtime 用户，使用 muModuleLoad 的 MUSA Driver 用户不受影响。
2. 内核加载优化：调用 muModuleLoad*() 时不会立即加载内核，而是延迟到调用 muModuleGetFunction() 时才加载。例外：某些内核（如指针存储在全局变量中的内核）仍需在调用 muModuleLoad*() 时加载。此优化同时影响 MUSA Runtime 和 MUSA Driver 用户。MUSA Runtime 仅在首次使用/引用内核时调用 muModuleGetFunction()。

这些优化对用户透明，开发者只要遵循标准的 MUSA 开发规范，即可实现底层性能优化特性的无缝衔接与自动使能。

2.4.1 延迟模式下的内核触发加载

内核和变量的加载是自动进行的，无须显式操作。只需启动内核或引用变量/内核，系统就会自动加载相关模块和内核。但如果需要无执行或无修改地预加载内核，建议采用以下方式：

• MUSA Driver API：内核加载通过调用 muModuleGetFunction() 触发。该调用在非延迟加载场景中同样必要，因为这是获取内核句柄的唯一途径。MUSA Driver API 还可实现更精细的加载控制。
• MUSA Runtime API：由于 MUSA Runtime API 自动管理模块，建议通过 musaFuncGetAttributes() 引用来触发内核加载。这种方式可确保加载内核时不改变其状态。

2.4.2 延迟加载状态查询

使用 MUresult muModuleGetLoadingMode(MUmoduleLoadingMode* mode) 可检测延迟加载是否启用。注意调用前需初始化 MUSA：

#include "musa.h"
#include "assert.h"
#include "iostream"
int main() {
    MUmoduleLoadingMode mode;
    assert(MUSA_SUCCESS == muInit(0));
    assert(MUSA_SUCCESS == muModuleGetLoadingMode(&mode));
    std::cout << "MUSA Module Loading Mode is " << ((mode == MU_MODULE_LAZY_LOADING) ?
    "lazy" : "eager") << std::endl;
    return 0;
}

2.4.3 采用延迟加载的潜在问题

延迟加载设计上无须修改应用代码，但若应用未完全遵循 MUSA 编程模型，可能出现以下三类问题：

三、虚拟内存管理：Driver API 层面的进阶能力

MUSA 虚拟内存管理是 Driver API 层面的进阶能力——核心是"直接管控 GPU 虚拟地址空间，并按需分配物理显存"，支持多进程/多 GPU 内存共享、自定义内存池、稀疏数据结构等，以满足 AI 训练中的数据共享、科学计算中的稀疏矩阵处理等需求。

3.1 核心 API 函数一览

3.2 Step-by-Step：预留虚拟地址 → 创建物理显存 → 映射 → 设置访问权限

1. 分配物理显存：使用 MUmemGenericAllocationHandle 表示物理显存，创建前需配置其属性（如分配位置、物理属性），并保证分配大小与显存粒度对齐（粒度可用 muMemGetAllocationGranularity() 查询）。

   size_t granularity = 0;
   MUmemAllocationProp prop = {};
   prop.type = MU_MEM_ALLOCATION_TYPE_PINNED;
   prop.location.type = MU_MEM_LOCATION_TYPE_DEVICE;
   prop.location.id = device;
   muMemGetAllocationGranularity(&granularity, &prop, MU_MEM_ALLOC_GRANULARITY_MINIMUM);
   alloc_size = ((alloc_size + granularity - 1) / granularity) * granularity;
   MUmemGenericAllocationHandle handle;
   muMemCreate(&handle, alloc_size, &prop, 0);

2. 预留虚拟地址：muMemAddressReserve() 为应用程序保留一块连续的虚拟地址空间，此时该区域并未绑定任何实际的物理显存资源，可根据需要映射到来自系统中任意 GPU 设备的物理显存块之上。

   MUdeviceptr va_ptr;
   muMemAddressReserve(&va_ptr, alloc_size, 0, 0, 0);

3. 建立映射：使用 muMemMap() 建立从虚拟地址到已分配物理内存的映射。

   muMemMap(va_ptr, alloc_size, 0, handle, 0);

4. 设置访问属性：映射后该地址还不能直接从设备访问，必须使用 muMemSetAccess() 设置访问属性。

   MUmemAccessDesc accessDesc = {};
   accessDesc.location.type = MU_MEM_LOCATION_TYPE_DEVICE;
   accessDesc.location.id = device;
   accessDesc.flags = MU_MEM_ACCESS_FLAGS_PROT_READWRITE;
   muMemSetAccess(va_ptr, alloc_size, &accessDesc, 1);

5. 资源清理（顺序与创建相反）：先取消映射，再释放虚拟地址，最后释放物理显存。

   muMemUnmap(va_ptr, alloc_size);
   muMemAddressFree(va_ptr, alloc_size);
   muMemRelease(handle);

具体示例请查看原书第 6 章代码文件 musaVirtualMemoryManagement.cpp。

3.3 高级特性与应用场景

当前 MUSA 虚拟内存管理已实现单机内多进程、多 GPU 间的内存共享，而未来将进一步支持 Fabric Handle 类型的虚拟内存创建，突破节点边界限制，使分配的物理显存可通过高速互联网络（如 InfiniBand、MTLink）实现跨节点共享，无须依赖数据复制或中间存储转发。该特性将与 MCCL 形成协同，进一步完善多节点并行计算的软件支撑。

• 内存共享：适用于多进程协作或多 GPU 协同场景（如分布式推理、多卡数据预处理）。通过 muMemExportToShareableHandle() 与 muMemImportFromShareableHandle() 函数，物理显存可在不同进程、不同 GPU 间直接共享，无须复制数据。
• 内存池：通过"预留虚拟地址 + 按需映射物理显存"的组合，构建自定义内存池，避免频繁调用 musaMalloc()/musaFree() 导致的开销，同时解决内存碎片化问题。
• 稀疏数据结构：适配科学计算中的稀疏矩阵、大规模点云数据等场景，仅为有效数据分配物理显存并建立虚拟地址映射，无效数据不占用硬件资源。

四、纹理内存：面向空间访问模式的专用内存

纹理内存是 MUSA 中面向特定场景深度优化的重要内存类型，在图形渲染与通用计算领域发挥关键作用，也可作为跨 API 数据共享的核心载体。

4.1 纹理内存核心特点

• 只读：纹理内存是只读的，不能直接写入。
• 缓存机制：专用的纹理缓存（L1 缓存）优化空间局部性。
• 边界处理：自动处理越界访问。
• 插值功能：支持硬件级线性插值。
• 格式转换：支持不同数据类型的自动转换。
• 访问方式：通过专门的纹理采样函数（tex1D、tex2D 等）访问。

适用场景：具有空间局部性的访问模式（如图像处理）；需要坐标变换（如图像旋转、缩放）的算法；需要边界条件处理的应用；需要硬件插值的数据处理；非对齐/非合并的内存访问模式。

4.2 Step-by-Step：纹理对象全生命周期

本节按"MUSA 数组创建 → 纹理参数配置 → 纹理对象创建 → 核函数调用 → 资源清理"的顺序拆解完整流程。

1. 创建 MUSA 数组：纹理内存的数据需依托 MUSA 数组存储，这是后续纹理参数配置与纹理对象创建的前提。

   // 定义通道格式描述符（32 位浮点数）
   musaChannelFormatDesc channelDesc = musaCreateChannelDesc(32, 0, 0, 0, musaChannelFormatKindFloat);
   // 创建二维数组
   musaArray* muArray;
   musaMallocArray(&muArray, &channelDesc, width, height);
   // 将主机数据复制到数组
   musaMemcpy2DToArray(
     muArray,                    // 目标数组
     0,                          // 目标 x 偏移
     0,                          // 目标 y 偏移
     hostData,                   // 主机数据源
     srcPitch,                   // 主机数据行间距
     width * sizeof(float),      // 实际复制宽度
     height,                     // 高度
     musaMemcpyHostToDevice);

2. 配置纹理参数：以图像平滑缩放为例，调整寻址模式、过滤模式等核心参数，参数直接决定访问性能与功能实现效果。

   // 资源描述符
   musaResourceDesc resDesc;
   memset(&resDesc, 0, sizeof(resDesc));                    // 清零初始化
   resDesc.resType = musaResourceTypeArray;                 // 资源类型为数组
   resDesc.res.array.array = muArray;                       // 指定绑定的 MUSA 数组
   
   // 纹理描述符
   musaTextureDesc texDesc;
   memset(&texDesc, 0, sizeof(texDesc));
   texDesc.addressMode[0] = musaAddressModeClamp;  // x 方向寻址模式：边界截断
   texDesc.addressMode[1] = musaAddressModeClamp;  // y 方向寻址模式：边界截断
   texDesc.filterMode     = musaFilterModeLinear;  // 过滤模式：线性插值
   texDesc.readMode       = musaReadModeElementType; // 读取模式：返回原始数据类型
   texDesc.normalizedCoords = 1;                   // 使用归一化坐标 [0,1]
   texDesc.borderColor[0]  = 0.0f;                 // 边界颜色（RGBA）
   texDesc.borderColor[1]  = 0.0f;
   texDesc.borderColor[2]  = 0.0f;
   texDesc.borderColor[3]  = 0.0f;

3. 创建纹理对象：纹理对象是核函数访问纹理内存的桥梁。

   musaTextureObject_t texObj = 0;
   musaError_t err = musaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
   if (err != musaSuccess)
   {
     printf("纹理对象创建失败: %s
   ", musaGetErrorString(err));
     // 错误处理
   }

4. 核函数调用：通过归一化坐标访问纹理，结合线性插值模式实现图像的平滑缩放。

   __global__ void kernel(musaTextureObject_t texObj, float* output, int width, int height) {
       int x = blockIdx.x * blockDim.x + threadIdx.x;
       int y = blockIdx.y * blockDim.y + threadIdx.y;
   
       if (x < width && y < height) {
         // 归一化坐标
         float u = (x + 0.5f) / width;
         float v = (y + 0.5f) / height;
         // 纹理采样（含硬件插值）
         float value = tex2D(texObj, u, v);
         output[y * width + x] = value;
       }
   }

5. 资源清理：

   // 销毁纹理对象
   musaDestroyTextureObject(texObj);
   // 释放 MUSA 数组
   musaFreeArray(muArray);

4.3 核函数中的纹理访问：一维/二维/三维

一维纹理适用于序列数据（如时间序列、一维信号），支持归一化与非归一化两种坐标模式：

// 非归一化坐标直接访问（无插值）
float value = tex1Dfetch(texObj, idx);
// 归一化坐标访问（带插值）
float u = (idx + 0.5f) / size;
float value = tex1D(texObj, u);

二维纹理是最常用的类型，广泛应用于图像处理场景。以下代码结合边界检查与中心采样，实现图像反色处理：

__global__ void processImage(musaTextureObject_t texObj, uchar4* output, int width, int height)
{
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;

  // 边界检查
  if (x >= width || y >= height) return;

  // 归一化坐标（使用中心采样避免偏移）
  float u = (x + 0.5f) / width;
  float v = (y + 0.5f) / height;

  // 带边界处理的纹理获取
  uchar4 pixel = tex2D(texObj, u, v);

  // 后处理
  pixel.x = 255 - pixel.x; // 反色处理示例
  output[y * width + x] = pixel;
}

三维纹理适用于立体数据（如三维模型、医学影像），访问逻辑与二维纹理类似，但需增加深度方向的归一化坐标：

float value = tex3D(texObj, u, v, w);

三维纹理的访问充分利用了纹理缓存的空间局部性，相比全局内存访问，能显著提高立体数据的读取效率。

4.4 纹理参数详解

归一化坐标示例：无符号 8 位整数 0→0.0，255→1.0；有符号 16 位整数 -32768→-1.0，32767→1.0。归一化坐标下的中心采样写法：

float u = (x + 0.5f) / width;

图像旋转是整合纹理内存"坐标变换 + 硬件插值"核心能力的典型应用场景，具体示例请查看原书第 6 章代码文件 imageRotation.cpp。

4.5 常见问题排查

• 纹理绑定失败：检查 MUSA 数组分配是否正确；验证资源描述符的 resType 设置；检查通道格式描述符与数据类型是否匹配。
• 读取结果不正确：检查归一化坐标是否正确计算；验证寻址模式和边界处理配置；确认读模式与纹理数据类型是否匹配。
• 性能不如预期：检查访问模式的空间局部性；尝试用不同块大小优化数据局部性；考虑使用常量内存存储纹理参数。
• 纹理尺寸限制：一维纹理最大尺寸 131072（计算能力 2.0+）；二维纹理最大尺寸 65536×32768（计算能力 3.0+）；三维纹理最大尺寸 4096×4096×4096。
• 多 GPU 环境适配：每个设备需独立绑定纹理；使用统一寻址可简化多 GPU 纹理访问；避免设备间同步纹理操作。

4.6 其他类型的纹理：表面内存、分层纹理、立方体贴图

表面内存突破基础纹理只读限制，支持设备端写入，适用于动态纹理更新（如实时渲染、视频帧处理场景）：

// 创建表面对象
musaSurfaceObject_t surfObj;
// 表面写入
__device__ void writeToSurface(musaSurfaceObject_t surf, int x, int y, uchar4 value) {
    surf2Dwrite(value, surf, x * sizeof(uchar4), y);
}
// 表面读取
__device__ uchar4 readFromSurface(musaSurfaceObject_t surf, int x, int y) {
    return surf2Dread(surf, x * sizeof(uchar4), y);
}

分层纹理本质是多个同尺寸纹理的集合，支持批量管理与访问，适用于多图层图像处理（如批量图像分类、多层纹理映射）场景：

musaArray* layeredArray;
musaExtent extent = make_musaExtent(width, height, layers);
musaMalloc3DArray(&layeredArray, &channelDesc, extent, musaArrayLayered);
// 访问方式
__device__ float readLayer(musaTextureObject_t tex, int layer, float u, float v) {
    return tex2DLayered(tex, u, v, layer);
}

立方体贴图由 6 个正方形纹理组成（对应立方体的 6 个面），适用于三维模型环境映射场景：

musaArray* cubemapArray;
musaMallocArray(&cubemapArray, &channelDesc, width, height, 6, musaArrayCubemap);
// 访问方式
__device__ float4 readCubeMap(musaTextureObject_t tex, float x, float y, float z) {
    return texCubemap(tex, x, y, z);
}

五、异步执行模型：流与事件

异步执行模型是 MUSA 并行计算的核心优化手段，其核心目标是打破主机与设备、设备内部任务间的执行阻塞，通过灵活的任务调度与精确的同步控制，实现计算、数据传输等操作的并行执行。该模型依托两大关键部分构建：流作为异步操作的组织载体，负责将一系列 MUSA 操作按序编排并支持多流间并行执行；事件作为流中的"标记点"，发挥同步控制与性能监控的核心作用。

5.1 流管理

流是一组按顺序执行的 MUSA 操作，这些操作在 GPU 上异步执行。主机代码可以在流执行期间继续运行，实现主机与设备的并行工作。流的核心特性：同一流内的操作按顺序执行；不同流之间的操作可以并行执行；支持数据传输与核函数计算的重叠；提供细粒度的同步控制。

Step-by-Step：流的创建、使用与销毁

1. 创建流：

   #include <musa_runtime.h>
   musaStream_t stream;
   musaStreamCreate(&stream);

2. 将异步操作关联到流：所有 MUSA 操作均可通过 stream 参数指定关联的流，以实现异步执行。

   // 异步内存操作
   musaMemcpyAsync(dst, src, size, direction, stream);
   // 异步核函数调用
   kernel<<<grid, block, shared_mem, stream>>>(args...);
   // 流回调（主机函数在流中指定点执行）
   musaStreamAddCallback(stream, callback_func, user_data, 0);

3. 按需同步：流提供多种同步方式，满足不同场景的顺序要求。

   // 1. 等待特定流完成所有操作
   musaStreamSynchronize(stream);
   
   // 2. 检查流是否完成所有操作（非阻塞）
   musaError_t status = musaStreamQuery(stream);
   if (status == musaSuccess) { /* 流已完成 */ }
   
   // 3. 事件同步：在流中插入事件并等待
   musaEvent_t event;
   musaEventCreate(&event);
   musaEventRecord(event, stream);
   musaEventSynchronize(event);
   
   // 4. 设备级同步（谨慎使用）
   musaDeviceSynchronize();

4. 销毁流（确保所有操作完成后再销毁）：

   musaStreamSynchronize(stream);
   musaStreamDestroy(stream);

流间并行与依赖控制

// 两个流并行执行不同任务
musaMemcpyAsync(d_a, h_a, size, musaMemcpyHostToDevice, stream1);
kernel1<<<grid, block, 0, stream1>>>(d_a);
musaMemcpyAsync(h_result1, d_result1, size, musaMemcpyDeviceToHost, stream1);

musaMemcpyAsync(d_b, h_b, size, musaMemcpyHostToDevice, stream2);
kernel2<<<grid, block, 0, stream2>>>(d_b);
musaMemcpyAsync(h_result2, d_result2, size, musaMemcpyDeviceToHost, stream2);

流与内存操作优化：计算与传输重叠

// 高效模式：计算与数据传输重叠
for (int i = 0; i < num_chunks; i++) {
    // 异步传输第 i 块数据
    musaMemcpyAsync(d_data + i*chunk_size,
                     h_data + i*chunk_size,
                     chunk_size,
                     musaMemcpyHostToDevice,
                     stream);
    // 异步处理第 i 块数据（与下一块数据的传输重叠）
    kernel<<<grid, block, 0, stream>>>(d_data + i*chunk_size, ...);
    // 异步回传结果
    musaMemcpyAsync(h_result + i*chunk_size,
                     d_result + i*chunk_size,
                     chunk_size,
                     musaMemcpyDeviceToHost,
                     stream);
}

默认流行为：所有未指定流的操作都在默认流（NULL 流）中执行；默认流与所有其他流同步（可能导致性能瓶颈）。

musaMemcpy(dst, src, size, direction);   // 无 stream 参数 → 默认流
kernel<<<grid, block>>>(args...);        // 默认流

流优先级：对于多流场景，可通过设置流优先级控制资源分配，让关键任务获得更多资源。

// 获取当前设备支持的流优先级范围
int priority_high, priority_low;
musaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

// 创建高优先级流和低优先级流
musaStream_t high_priority_stream, low_priority_stream;
musaStreamCreateWithPriority(&high_priority_stream, musaStreamNonBlocking, priority_high);
musaStreamCreateWithPriority(&low_priority_stream, musaStreamNonBlocking, priority_low);

5.2 事件管理

如果说流是"异步任务的组织者"，事件就是"异步任务的控制器"。事件是 GPU 编程中用于同步操作和精确计时的核心机制，作用包括：监控 GPU 任务进度、测量操作时间间隔、实现跨流依赖控制、确保执行顺序的正确性。

Step-by-Step：事件的创建、记录、同步、销毁

1. 创建事件对象：

   musaEvent_t event;
   musaEventCreate(&event);
   // 创建带标志的事件（如禁用计时）
   musaEventCreateWithFlags(&event, musaEventDisableTiming);

   标志选项说明：musaEventDefault（默认设置）、musaEventBlockingSync（主机同步使用阻塞模式）、musaEventDisableTiming（禁用计时功能，通常在不需要使用事件的计时功能时设置）、musaEventInterprocess（支持进程间共享，设置后可能会影响性能，应仅在需要进行进程间通信的场景下设置）。

2. 在流中记录事件并同步：

   // 在指定流中记录事件（默认为 0 流）
   musaEventRecord(event, stream);
   // 阻塞主机直到事件完成
   musaEventSynchronize(event);
   // 非阻塞查询（返回 musaSuccess 或 musaErrorNotReady）
   musaError_t status = musaEventQuery(event);

   两种同步方式的核心区别：musaEventSynchronize 是"主动等待"，会阻塞主机直到事件完成；musaEventQuery 是"被动查询"，主机可继续执行其他任务，适用于异步轮询场景。

3. 测量时间间隔：

   float elapsedTime;
   musaEventElapsedTime(&elapsedTime, startEvent, stopEvent); // 单位: ms

   测量时间间隔时，分辨率通常为 0.5μs。startEvent 和 stopEvent 必须在同一上下文，且事件创建时未设置 musaEventDisableTiming。

4. 销毁事件（避免资源泄漏）：

   musaEventDestroy(event);

典型应用场景一：内核执行时间测量

musaEvent_t start, stop;
musaEventCreate(&start);
musaEventCreate(&stop);
musaEventRecord(start);   // 记录起始点
myKernel<<<grid, block>>>(...);
musaEventRecord(stop);    // 记录结束点
musaEventSynchronize(stop);
float time;
musaEventElapsedTime(&time, start, stop);
printf("Kernel time: %.3f ms
", time);
// 清理资源
musaEventDestroy(start);
musaEventDestroy(stop);

典型应用场景二：跨流同步

当流 B 的任务依赖流 A 时，可通过事件实现非阻塞同步，避免主机介入：

musaStream_t stream1, stream2;
musaEvent_t syncEvent;
musaStreamCreate(&stream1);
musaStreamCreate(&stream2);
musaEventCreate(&syncEvent);
// 在 stream1 中记录事件
kernelA<<<..., stream1>>>(...);
musaEventRecord(syncEvent, stream1);
// stream2 等待事件完成
musaStreamWaitEvent(stream2, syncEvent, 0);
kernelB<<<..., stream2>>>(...);

关键机制：musaStreamWaitEvent 用于建立非阻塞流之间的依赖，确保 kernelB 在 kernelA 后执行，无须在主机端同步；该函数还支持传入来自与当前流处于不同上下文和设备的事件。

典型应用场景三：重叠计算与数据传输

musaStream_t computeStream, dataStream;
musaEvent_t dataReady1, dataReady2;
// 异步传输数据
musaMemcpyAsync(devData1, hostData1, size, musaMemcpyHostToDevice, dataStream);
musaEventRecord(dataReady1, dataStream);
musaMemcpyAsync(devData2, hostData2, size, musaMemcpyHostToDevice, dataStream);
musaEventRecord(dataReady2, dataStream);
// 计算流等待数据就绪
musaStreamWaitEvent(computeStream, dataReady1, 0);
kernel<<<..., computeStream>>>(devData1);
musaStreamWaitEvent(computeStream, dataReady2, 0);
kernel<<<..., computeStream>>>(devData2);

性能最佳实践与常见问题

事件复用：频繁创建/销毁事件会产生额外开销，建议复用事件。

// 初始化时创建
static musaEvent_t timerStart, timerStop;
musaEventCreate(&timerStart);
musaEventCreate(&timerStop);
// 每次测量重用
musaEventRecord(timerStart);
kernel<<<...>>>();
musaEventRecord(timerStop);

异步事件处理：对于耗时较长的 GPU 任务，可通过非阻塞查询事件状态，让主机并行处理其他任务，适用于 GPU 任务耗时远大于 CPU 任务耗时的场景。

while (musaEventQuery(event) == musaErrorNotReady) {
    do_cpu_work();  // 并行执行 CPU 任务
}

六、同步机制：保证并行计算结果正确性

MUSA 编程模型中，主机端由用户负责管理和调度任务，设备端则根据依赖关系处理大规模并行计算任务。在并行计算过程中，不同线程可能会同时访问和修改共享数据，若不加以管理，就会出现数据竞争和数据不一致的问题。同步机制的作用就是协调这些并行计算线程的执行顺序，确保在特定操作完成之前，其他线程不会访问或修改相关数据，从而保证计算结果的正确性。

6.1 线程级同步：__syncthreads()

__syncthreads() 函数是线程块内同步的核心工具。它会阻塞线程块内所有线程的执行，直到块内的每个线程都执行到该函数调用处，然后所有线程再继续向下执行，从而确保线程块内的所有线程在某个操作点达到同步状态，保证数据的一致性。

__global__ void blockSyncKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = threadIdx.x;

    // 线程加载全局内存数据到共享存储器
    if (tid < n) {
        sharedData[tid] = data[tid];
    }
    __syncthreads();
    // 所有线程都完成将数据加载到线程间共享存储器的操作后，再进行后续操作
    if (tid < n - 1) {
        sharedData[tid] += sharedData[tid + 1];
    }
}

在以上示例代码中，线程首先将数据从全局内存加载到线程间共享存储器，通过 __syncthreads() 确保所有线程都完成加载后，再进行数据累加操作。

6.2 任务级同步：musaDeviceSynchronize() 与 musaStreamSynchronize()

任务级同步主要针对主机与设备、流与流之间的协同，其核心目标是确保"依赖任务按顺序执行"。

musaDeviceSynchronize()：主机与设备之间的同步

该函数会阻塞主机线程的执行，直到设备上所有已提交的内核函数都执行完毕。调用后，主机可以安全地访问设备计算后的结果，而不用担心数据还未计算完成。

float* h_data = new float[N];
float* d_data;
musaMalloc((void**)&d_data, N * sizeof(float));
musaMemcpy(d_data, h_data, N * sizeof(float), musaMemcpyHostToDevice);
kernel<<<gridSize, blockSize>>>(d_data, N);
musaDeviceSynchronize();
musaMemcpy(h_data, d_data, N * sizeof(float), musaMemcpyDeviceToHost);
musaFree(d_data);
delete[] h_data;

musaStreamSynchronize()：同步指定流

该函数用于同步指定的流，它会阻塞主机线程，直到指定流中的所有操作都已完成。

musaStream_t stream1, stream2;
musaStreamCreate(&stream1);
musaStreamCreate(&stream2);
// 在 stream1 中执行操作
musaMemcpyAsync(d_data1, h_data1, N * sizeof(float), musaMemcpyHostToDevice, stream1);
kernel1<<<gridSize, blockSize, 0, stream1>>>(d_data1, N);
// 在 stream2 中执行操作
musaMemcpyAsync(d_data2, h_data2, N * sizeof(float), musaMemcpyHostToDevice, stream2);
kernel2<<<gridSize, blockSize, 0, stream2>>>(d_data2, N);
// 同步 stream1
musaStreamSynchronize(stream1);
// 此时可以安全地访问 stream1 计算后的结果

通过使用 musaStreamSynchronize()，开发者可以更精细地控制流内操作的完成时间。任务级的同步在需要确保整个计算任务完成后再进行后续处理的场景（如深度学习训练）中至关重要；在数据处理管道中，当多个 MUSA 内核函数依次处理数据时，也需要同步来保证数据处理顺序的正确性。

6.3 MUSA API 的默认同步行为

例如，musaMemcpyAsync(d_data, h_data, N * sizeof(float), musaMemcpyHostToDevice, stream) 通过指定 MUSA stream，将数据从主机内存 h_data 异步复制到设备内存 d_data；主机在调用该函数后可立即继续执行后续代码，从而在数据传输的同时进行其他不相关计算或任务调度，充分利用 CPU 资源。

七、MUSA Graph 管理：批量提交与执行优化

MUSA Graph 是摩尔线程提供的一种机制，用于捕获和重放一系列 MUSA 操作，以优化应用程序的性能。传统编程模式中，每次提交内核函数、进行异步内存复制等操作时，CPU 都需要与 GPU 驱动进行多次交互，这会产生大量启动开销，还会因 CPU 频繁参与调度而限制整体执行效率。MUSA Graph 通过将一组有依赖关系的 MUSA 操作封装为一个可复用的"计算图"，实现操作序列的批量提交与优化执行。

7.1 三个核心概念

• MUSA Graph：表示一系列 MUSA 操作的依赖关系图，本质是包含多个操作节点的有向无环图，明确记录节点间的执行顺序（比如"数据复制完成后启动核函数"）。
• MUSA Node：图中的节点，代表一个独立操作。常见节点类型包括内核启动节点、内存复制节点（对应 musaMemcpyAsync 类操作）、事件节点（用于同步）等。
• Graph Instance：可执行的图实例。Graph 仅为操作的逻辑定义，需通过实例化转换为可在 GPU 上运行的执行计划，实例化过程中驱动会对图内操作进行底层优化（如调整执行顺序、合并冗余操作）。一个 Graph 可生成多个独立的 Instance。

7.2 三大核心优势

• 降低启动开销：Graph 将多个操作的启动流程合并为一次，避免重复启动开销，在高频重复执行相同操作序列时优势会被显著放大。
• 优化执行顺序：驱动和运行时库可根据 Graph 内的依赖关系对操作执行顺序进行智能重排（前提是不破坏依赖逻辑），例如将可并行的内存复制与核函数操作调度到不同流中同时执行。
• 减少 CPU 参与：Graph 被捕获和实例化后，只需 CPU 触发一次 Graph 启动指令，CPU 可转向处理其他任务。

7.3 核心 API 与两种构建方式

musaGraphCreate()           // 创建空图
musaGraphDestroy()          // 销毁图
musaGraphAddKernelNode()    // 添加内核节点
musaGraphAddMemcpyNode()    // 添加内存复制节点
musaStreamBeginCapture()    // 开始捕获流中的操作
musaStreamEndCapture()      // 结束捕获并返回图
musaGraphInstantiate()      // 实例化图
musaGraphLaunch()           // 启动图执行

方式一：显式构建（精准控制依赖关系）

musaGraph_t graph;
musaGraphCreate(&graph, 0);
musaGraphNode_t kernelNode;
musaKernelNodeParams params = {…};
musaGraphAddKernelNode(&kernelNode, graph, NULL, 0, ¶ms);
musaGraphExec_t instance;
musaGraphInstantiate(&instance, graph, NULL, NULL, 0);
musaGraphLaunch(instance, stream);

显式构建的优势在于灵活性高，开发者可按需为每个节点指定依赖关系，适用于处理对执行逻辑有严格要求的复杂任务，例如包含多阶段核函数、多轮数据复制的流水线任务。

方式二：流捕获（零侵入式性能优化）

musaStreamBeginCapture(stream, musaStreamCaptureModeGlobal);
kernel<<<…, stream>>>(…);
musaMemcpyAsync(…, stream);
musaGraph_t graph;
musaStreamEndCapture(stream, &graph);
musaGraphExec_t instance;
musaGraphInstantiate(&instance, graph, NULL, NULL, 0);
musaGraphLaunch(instance, stream);

流捕获的优势在于接入成本低，无须修改现有操作的提交逻辑，只需在操作序列前后添加捕获启停指令，适用于对现有项目进行轻量化性能优化。

7.4 高级特性

• 更新机制：允许更新图实例而不重新实例化图。部分操作参数可能动态变化（如核函数的输入数据指针、网格维度），更新机制可直接修改实例中的指定参数，避免重复实例化的额外开销。
• 事件记录：支持在图中添加事件记录节点，标记关键操作的执行起止时间，方便后续通过性能分析工具定位执行瓶颈，也可作为同步节点实现图内部或图之间的精准同步。
• 条件节点：支持条件执行分支，可根据指定条件（如某一变量的取值）选择执行不同的操作分支，使图具备处理分支逻辑的能力。
• 子图：支持图的嵌套，适用于构建层次化的复杂计算流程。可将一组高频复用的操作序列封装为子图，再将子图作为一个节点挂载到主图中。

八、MUSA 图形互操作：与 OpenGL 等图形 API 协同

MUSA 图形互操作（MUSA Graphics Interoperability）旨在实现 MUSA 并行计算与图形 API（如 OpenGL、Direct3D、Vulkan）间的低开销数据共享。其核心价值在于避免 GPU 内存的冗余复制，通过资源共享机制，提升实时图形处理、科学可视化以及深度学习等高性能应用的运行效率。

整个过程包括资源注册和资源映射管理两大环节：

• 资源必须通过 MUSA interOp() 函数注册至 MUSA，注册函数会返回 struct musaGraphicsResource 类型的指针。由于注册操作存在较高开销，通常每个资源仅需注册一次，且每个需要使用资源的 MUSA 上下文必须单独注册。不再使用时可通过 musaGraphicsUnregisterResource() 注销。
• 资源注册后可通过 musaGraphicsMapResources()、musaGraphicsUnmapResources() 函数进行映射管理，同一资源支持反复映射与解除映射操作。还可通过 musaGraphicsResourceSetMapFlags() 设置使用提示（如只读、只写）以进一步提升资源访问效率。
• Kernel 可以通过 musaGraphicsResourceGetMappedPointer() 获得缓冲区设备内存地址；对于 MUSA 数组，也可通过 musaGraphicsSubResourceGetMappedArray() 获得句柄。

8.1 OpenGL 互操作

OpenGL 中允许映射至 MUSA 地址空间的资源包括缓冲区对象（Buffer Object）、纹理对象以及渲染缓冲区对象（Renderbuffer Object）三类。

缓冲区对象可以通过 musaGraphicsGLRegisterBuffer() 完成注册。注册后，该对象在 MUSA 中表现为线性设备内存指针，支持标准内存操作；Kernel 可直接对注册后获得的设备指针进行读写，同时该指针也支持 musaMemcpy() 操作。

纹理对象和渲染缓冲区对象可以通过 musaGraphicsGLRegisterImage() 进行注册。注册后获得的 MUSA 数组需要绑定至纹理引用或表面引用（Surface Reference）进行读取，也可以通过 musaMemcpy2D() 进行读写操作。若在注册时启用 musaGraphicsRegisterFlagsSurfaceLoadStore 标志，则可通过 surf2Dwrite() 等函数直接写入数据。

musaGraphicsGLRegisterImage() 支持以下 3 种数据格式：浮点型（通常为 1/2/4 通道，常见格式为 GL_RGBA_FLOAT32）；规格化整数（常见格式为 GL_RGBA8、GL_INTENSITY16）；非规格化整数（仅 OpenGL 3.0 及以上版本支持，仅可通过着色器写入，固定功能管线不可用，常见格式为 GL_RGBA8UI）。

调用 MUSA-OpenGL 互操作函数（如 musaGraphicsGLRegister() 函数或资源映射函数）时，必须通过 glXMakeCurrent（Linux）确保目标 OpenGL 上下文已绑定至当前主机线程。

8.2 完整互操作示例代码

// 1. OpenGL 初始化
GLuint texID;
glGenTextures(1, &texID);
glBindTexture(GL_TEXTURE_2D, texID);
glTexImage2D(…, GL_LUMINANCE, …);   // 灰度格式

// 2. MUSA 资源注册
musaGraphicsResource* musaTexRes;
musaGraphicsGLRegisterImage(&musaTexRes, texID, GL_TEXTURE_2D,
            musaGraphicsRegisterFlagsSurfaceLoadStore);

// 3. 渲染循环
while(running) {
    // MUSA 端映射
    musaGraphicsMapResources(1, &musaTexRes, 0);
    musaArray* muArray;
    musaGraphicsSubResourceGetMappedArray(&muArray, musaTexRes, 0, 0);
    // 获取 Surface 对象
    musaSurfaceObject_t surfObj;
    musaResourceDesc desc = {.resType = musaResourceTypeArray, .res.array.array = muArray};
    musaCreateSurfaceObject(&surfObj, &desc);
    // MUSA 内核更新纹理
    kernel<<<blocks, threads>>>(surfObj, width, height);
    // 解除映射
    musaDestroySurfaceObject(surfObj);
    musaGraphicsUnmapResources(1, &musaTexRes, 0);
    // OpenGL 渲染
    glDrawElements(…);
}

版本与兼容性

本课小结

课后练习

1. 写出三种查询/管理 MUSA 设备的方式，并说明各自的适用场景（如何在多 GPU 环境下选出计算能力最强的设备）。
2. 简述 fatbin 文件的编译方式，说明 MUSA_MODULE_LOADING=LAZY 相对于 EAGER 的优势与潜在风险（至少各举一例）。
3. 结合本课虚拟内存管理 API 表，写出使用虚拟内存管理 API 完成"预留虚拟地址→创建物理显存→映射→设置访问权限"的完整代码骨架。
4. 纹理内存与全局内存相比有哪些核心特点？请说明纹理对象相较纹理引用的三个主要优势。
5. 参照本课流与事件示例，编写一段代码，使用两个流分别异步传输并处理两块数据，并用事件测量各自内核的执行时间。
6. 简述 musaDeviceSynchronize()、musaStreamSynchronize()、__syncthreads() 三者的作用范围与区别。
7. 简述 MUSA Graph 的两种构建方式（显式构建、流捕获）及各自适用场景，并说明为什么高频重复的推理任务适合使用 MUSA Graph。
8. 简述 MUSA-OpenGL 互操作的关键步骤（资源注册→映射→内核访问→解除映射），并说明为何"资源映射期间禁止通过 OpenGL 直接访问资源"。

与其他课程的衔接

• m11-l7（第 5 章）：本课 musaInfo 硬件参数字段的编程读取方式，正是对上一课环境搭建成果的延伸使用
• m11-l3（第 3 章）：本课的流、事件、同步机制建立在第 3 章的 Kernel/Context/Stream 编程模型之上
• m11-l4（第 4 章）：本课涉及的 muDNN 相关内存管理与第 4 章软件架构中的计算加速库分层相互印证
• m11-l9（第 7 章）：本课编写的 MUSA 程序，下一课将学习如何用 musa-gdb 与 compute-sanitizer 定位其中的错误
• m11-l5（第 8 章）：本课的流并行、Occupancy 相关设备属性字段，会在第 8 章性能优化中被系统量化