MUSA C++ 语法概述

为什么需要语法层

编程模型提供 GPU 并行计算的抽象概念：

• 线程层次结构：网格 → 线程块 → 线程
• 内存层次结构：全局/共享/寄存器
• 执行模型：SIMT 执行，线程束调度

MUSA C++ 语法将这些概念转换为可执行代码：

编程模型概念              语法表达                    API 调用
    ↓                       ↓                         ↓
线程层次结构  →   threadIdx, blockIdx    →   musaMalloc, musaMemcpy
内存层次结构  →   __shared__             →   musaMemPrefetchAsync
执行模型   →   __syncthreads()        →   musaDeviceSynchronize

详细语法参考：MUSA C++ 语言扩展

---

API 层次结构

MUSA 提供两层 API：

Runtime API vs Driver API 对比

特性	Runtime API	Driver API
初始化	自动（首次调用时）	手动（muInit()）
上下文管理	隐式（primary context）	显式（muCtxCreate()）
模块加载	自动（从 PTX/二进制）	手动（muModuleLoad()）
编程复杂度	低	高
灵活性	有限	完全控制
适用场景	应用程序、快速原型	框架、库、多租户

层次结构图

┌─────────────────────────────────────────────────┐
│           编程模型 (Programming Model)          │
│  Thread Hierarchy / Memory Hierarchy / ...      │
└─────────────────────────────────────────────────┘
                    ↓ 如何表达
┌─────────────────────────────────────────────────┐
│              语法层 (Syntax Layer)              │
│  __global__, __shared__, threadIdx, <<<>>>      │
└─────────────────────────────────────────────────┘
                    ↓ 如何调用
┌─────────────────────────────────────────────────┐
│            API 层 (API Layer)                   │
│  ┌─────────────────┐     ┌─────────────────┐    │
│  │  Runtime API    │     │   Driver API    │    │
│  │  (高层接口)     │     │   (底层接口)    │     │
│  │                 │     │                 │    │
│  │  musaMalloc     │     │  muMemAlloc     │    │
│  │  musaMemcpy     │     │  muMemcpy       │    │
│  │  musaSetDevice  │     │  muCtxCreate    │    │
│  └─────────────────┘     └─────────────────┘    │
└─────────────────────────────────────────────────┘
                    ↓ 如何使用
┌─────────────────────────────────────────────────┐
│            设备端函数 (Device Functions)        │
│  原子函数 / Warp 函数 / 数学函数                 │
└─────────────────────────────────────────────────┘

代码对比：同一操作的两种 API

分配内存并运行 Kernel：

// ═══════════════════════════════════════════════════════
// Runtime API（简洁）
// ═══════════════════════════════════════════════════════
float *d_data;
musaMalloc(&d_data, size);                    // 自动初始化
musaMemcpy(d_data, h_data, size, ...);        // 自动上下文
myKernel<<<grid, block>>>(d_data);            // 直接启动
musaDeviceSynchronize();                      // 自动等待

// ═══════════════════════════════════════════════════════
// Driver API（灵活）
// ═══════════════════════════════════════════════════════
muInit(0);                                    // 手动初始化
muDeviceGet(&device, 0);
muCtxCreate(&context, 0, device);            // 创建上下文
muModuleLoad(&module, "kernel.mubin");         // 加载模块
muModuleGetFunction(&kernel, module, ...);   // 获取函数
muMemAlloc(&d_data, size);                   // 分配内存
muMemcpy(d_data, h_data, size, ...);         // 拷贝数据
muLaunchKernel(kernel, grid, block, ...);    // 启动 Kernel
muCtxSynchronize(context);                   // 同步上下文
muCtxDestroy(context);                       // 清理上下文

---

使用路径

1. 语法入门：MUSA C++ 入门
2. API 选择：应用开发 → Runtime API 指南，框架开发 → Driver API 指南
3. 设备端函数：原子函数，Warp 函数
4. 完整参考：Driver API，Runtime API

---

相关文档

• MUSA C++ 语法详解：语言扩展说明
• Runtime API 指南：应用开发推荐
• Driver API 指南：框架开发推荐