第4章精讲：MUSA 软件架构（编译器工具链 / 计算加速库 / muDNN / MCCL）

一、MUSA 软件栈总览：五大组件分工

MUSA 软件栈允许开发者用 C/C++ 编写面向摩尔线程 GPU 的程序，并配套提供多领域加速库，服务嵌入式设备、数据中心、超级计算机等多种部署形态。其核心载体是 MUSA Toolkit，由五大组件构成：

可以把这五个组件类比成一条流水线：驱动/运行时负责"设备管得住"，编译器负责"代码编得出"，加速库负责"常用计算不用自己写"，AI/通信库负责"大模型、多卡跑得动"，性能工具负责"哪里慢看得清"。

二、驱动与运行时库：手动挡与自动挡

驱动与运行时库负责主机（CPU）与设备（GPU）之间的交互，提供两套互补的 API：

一句话记忆：Driver API 是"手动挡"，掌控精细但操作繁琐；Runtime API 是"自动挡"，省心但底层细节被隐藏。

互操作三条规则

1. 设备管理与版本管理可互换使用：两套 API 在这方面完全兼容。
2. 上下文创建的隐式/显式关系：Runtime API 被调用时会隐式自动创建上下文，称为主上下文（Primary Context）；若开发者先用 Driver API 主动创建并激活了一个上下文，后续 Runtime API 调用会复用该上下文而不会重新创建。Driver API 可以对主上下文进行完整管理。
3. 内存操作完全兼容，指针可双向转换：无论用 muMemAlloc() 还是 musaMalloc() 分配显存，得到的设备指针可以互相转换使用。

// Driver API 与 Runtime API 指针互转示例
MUdeviceptr devPtr;
float *d_data;

// 用 Driver API 分配
muMemAlloc(&devPtr, size);
d_data = (float *)devPtr;

// 用 Runtime API 分配
musaMalloc(&d_data, size);
devPtr = (MUdeviceptr)d_data;

这一互操作特性的现实意义是：用 Driver API 写的底层程序可以直接调用基于 Runtime API 构建的各类库，无需额外适配，这也是 muBLAS、muDNN 等高性能库能同时服务两类用户的原因。

三、编译器工具链：mcc 与 MTX

mcc 是 MUSA 编程语言配套的编译器工具链，负责编译构建 MUSA 应用程序。因为一份 .mu 源文件同时混杂主机端标准 C++ 代码与设备端扩展代码，mcc 需要把两条编译路径的复杂性隐藏起来，并针对不同 MT GPU 硬件架构做代码优化。

MTX 编译器（Moore Threads eXtension）是工具链中间层，专门处理 MTX 格式的中间表示代码，承接 mcc 输出的中间代码，完成向目标硬件指令的转换与优化。二者共同构成"高级语言 → 中间表示 → 机器码"的完整编译链路。

3.1 mcc 编译流程：从 .mu 到可执行文件

一个 .mu 源文件会被拆分成两条并行编译路径：

• 主机端路径：源码经 Clang++ 编译生成内存中的 LLVM IR，再经 llc（LLVM 静态编译器）生成 .o（主机目标文件）。
• 设备端路径：源码经 Clang++ 和 llc 生成 elf 格式的 .o（设备目标文件），再经 lld（LLVM 链接器）重定位生成 .out，最后经 clang-offload-bundler 打包成 .a（卸载捆绑包）。

两条路径最终在主机端目标文件层面汇合：主机目标文件 .o 与设备侧卸载捆绑包共同交给 ld（主机链接器）完成最终链接，生成可执行文件。图中还标注两条执行规则：每个 .mu 源码文件都要重复"实线框"内的流程；GPU 的每一种子架构都要重复"虚线框"内的流程，即同一份设备代码要为不同 GPU 子架构分别生成对应机器码。

mcc 支持的输入文件后缀：

mcc 的具体编译分三步：① 预处理与语言识别（区分主机端 C++ 代码与设备端扩展代码）；② 两次独立编译（分别生成主机代码对象与设备代码对象，都要经过"编译成 LLVM IR → 目标代码生成"子流程）；③ 链接整合生成可执行文件。典型编译命令：

mcc target.mu -lmusart -L/usr/local/musa/lib -O2 -o target

参数含义：target.mu 为输入源文件；-L<dir> 指定库搜索路径；-l<lib-name> 链接运行时库（如 musart）；-O2 表示优化级别；-o 指定输出文件名。

3.2 编译控制选项与内核属性

mcc 提供的命令行选项大多与 Clang/LLVM 一致，同时为 MUSA 语言和 MT GPU 目标扩展了专属选项：

• 前端选项：如 -Wno-shared-var-init，可忽略因初始化线程间共享存储器导致的编译错误。
• 后端选项：以 -mllvm 为前缀，例如 -mtgpu-maxregcnt=<value> 可设置 MT GPU 目标的最大临时寄存器计数，是性能调优的常用手段。

mcc 还支持通过 C++ __attribute__ 语法为 MT GPU 内核函数提供编译优化提示，常用属性：

// mtgpu_num_usreg：限制内核使用的 TEMP/ATTR 寄存器数量（0 表示无限制）
__attribute__((mtgpu_num_usreg(256))) void KernelName(...) { /* ... */ }

// mtgpu_tiny_offset：提示编译器该内核内存访问偏移量不超过 INT_MAX，
// 从而为编译器创造更多优化机会
__global__ __attribute__((mtgpu_tiny_offset)) void kernel_name(...) { /* ... */ }

// mtgpu_unroll_threshold：设置循环展开阈值，作为编译器优化提示
// 用法：__attribute__((mtgpu_unroll_threshold(N)))

这类属性本质是开发者向编译器"透露"关于寄存器规模、循环特征、内存访问范围的先验知识，属于性能调优的"高级选项"，在大规模内核、寄存器紧张场景中价值明显。

3.3 MTX 编译器：中间表示处理与目标代码生成枢纽

MTX 编译器承担三件事：

• 中间表示处理：接收 mcc 输出的 MTX 格式中间代码，该格式是专为 GPU 并行计算设计的指令集架构表示。
• 目标代码生成：将 MTX 中间代码转换为针对摩尔线程 GPU 优化的 LLVM IR，进而生成高效机器码。
• 编译优化支持：叠加针对 MT GPU 架构特性（MP/MPE 并行单元、分层存储体系、SIMD 执行模式）的专项优化，支持多级优化策略、离线编译模式与调试信息生成（可输出 LLVM IR 和汇编代码）。

MTX 指令集覆盖算术运算（含 FMA 乘加融合）、位操作、逻辑运算、比较与选择（支持谓词执行）、内存访问（覆盖全局/共享/常量等多种空间）、同步与通信、特殊函数（三角函数、指数函数等超越函数）等类别；数据类型支持标量整数（8/16/32/64 位）、浮点（FP16/FP32/FP64）、向量类型及谓词/地址类型。

小结记忆点：mcc 管"前端语法 + 两次编译 + 链接"，MTX 编译器管"中间表示 → 机器码的深度架构优化"，两者串联起完整编译链路。

四、计算加速库：muBLAS / muRAND / muFFT / muSPARSE / muSOLVER

4.1 muBLAS：注意列优先存储

muBLAS 提供 BLAS-1（标量/向量运算，如缩放、点积）、BLAS-2（矩阵-向量计算）、BLAS-3（矩阵-矩阵计算）三类算子，支持单/双精度实数复数及 FP16、BF16、FP8、INT8 等多种数据类型。基本调用流程为四步法：创建句柄 → 分配 GPU 存储 → 调用算子 → 释放资源。

#include "mublas.h"
#include 
#include 
#include 
#include 
#include 

// 矩阵维度: A (M x K), B (K x N), C (M x N)
#define DIM1 1024  // M
#define DIM2 1024  // N
#define DIM3 1024  // K

int main() {
    mublasOperation_t transa = MUBLAS_OP_N, transb = MUBLAS_OP_N;
    double alpha = 1.1, beta = 0.9;
    mublas_int m = DIM1, n = DIM2, k = DIM3;
    mublas_int lda, ldb, ldc, size_a, size_b, size_c;

    // muBLAS 为列优先存储（column-major），需根据是否转置计算主维跨度：
    // transa == MUBLAS_OP_N 时 lda = m，否则 lda = k
    // transb == MUBLAS_OP_N 时 ldb = k，否则 ldb = n
    // C 不转置：ldc = m
    // 此处按 transa/transb 分支具体赋值 lda/ldb/ldc/size_a/size_b/size_c（略，需按实际转置状态适配）

    std::vector ha(size_a), hb(size_b), hc(size_c);
    // 主机端随机初始化 ha/hb/hc（需自行补充初始化逻辑）

    double *da, *db, *dc;
    musaMalloc(&da, size_a * sizeof(double));
    musaMalloc(&db, size_b * sizeof(double));
    musaMalloc(&dc, size_c * sizeof(double));
    musaMemcpy(da, ha.data(), sizeof(double) * size_a, musaMemcpyHostToDevice);
    musaMemcpy(db, hb.data(), sizeof(double) * size_b, musaMemcpyHostToDevice);
    musaMemcpy(dc, hc.data(), sizeof(double) * size_c, musaMemcpyHostToDevice);

    mublasHandle_t handle;
    mublasCreate(&handle);

    // C = alpha * op(A) * op(B) + beta * C
    mublasDgemm(handle, transa, transb, m, n, k, &alpha, da, lda, db, ldb, &beta, dc, ldc);

    musaMemcpy(hc.data(), dc, sizeof(double) * size_c, musaMemcpyDeviceToHost);

    musaFree(da); musaFree(db); musaFree(dc);
    mublasDestroy(handle);
    return 0;
}

4.2 muRAND：主机端与设备端随机数生成

接口分主机端库（murand.h）与设备端头文件（murand_kernel.h）。实际随机数生成计算发生在 GPU 上，结果存储在设备全局内存中。支持 8 种生成器：伪随机数（XORWOW、MRG32K3A、MTGP32、Philox_4_32_10）与准随机数（Sobol32、Sobol64、Scrambled Sobol32、Scrambled Sobol64）。初始化可配置种子（seed）、偏移量（offset）、顺序（order）三个参数。

主机端生成流程（五步）：

// 伪代码骨架，需结合具体生成器类型与精度补充完整
murandGenerator_t gen;
murandCreateGenerator(&gen, MURAND_RNG_PSEUDO_XORWOW);      // ① 创建生成器
murandSetPseudoRandomGeneratorSeed(gen, 1234ULL);            // ② 设置种子/偏移/顺序
float *d_data;
musaMalloc((void**)&d_data, n * sizeof(float));              // ③ 设备端分配内存
murandGenerateUniform(gen, d_data, n);                        // ④ 生成随机数
murandDestroyGenerator(gen);                                  // ⑤ 释放资源

若在主机 CPU 上直接生成，第①步改用 murandCreateGeneratorHost()，第③步改为分配主机内存缓冲区。设备端核函数内可引用 murand_kernel.h 直接生成随机数：伪随机数需先调用 murand_init() 初始化状态，再调用 murand()、murand_uniform()、murand_normal()、murand_log_normal() 等；准随机数（Sobol 系列）初始化不依赖种子，而是配置方向向量（direction_vectors），额外支持 murand_poisson 生成泊松分布。

4.3 muFFT：注意 inplace 模式的显存对齐

muFFT 是纯 C 语言接口库，API 沿用主流傅里叶变换库的接口设计习惯（〔存疑〕原书未点名具体是哪一个"主流库"，只表述为"延续了主流傅里叶变换库的接口设计"，此处不做臆测认定）。

#include 
#include 
int main()
{
    const int Nx = 4096;
    std::vector hx(Nx);
    size_t complex_bytes = sizeof(decltype(hx)::value_type) * hx.size();

    mufftComplex *dx;
    musaMalloc(&dx, complex_bytes);
    musaMemcpy(dx, hx.data(), complex_bytes, musaMemcpyHostToDevice);

    mufftHandle plan = NULL;
    mufftPlan1d(&plan, Nx, MUFFT_C2C, 1);

    mufftExecC2C(plan, dx, dx, MUFFT_FORWARD);
    musaMemcpy(hx.data(), dx, complex_bytes, musaMemcpyDeviceToHost);

    mufftDestroy(plan);
    musaFree(dx);
    return 0;
}

流程要点：初始化数据 → 申请 GPU 资源并拷贝 → 创建执行计划（plan）→ 执行变换 → 释放计划和资源。执行计划仅对应一种类型和尺寸的变换，在释放之前可反复使用；最佳实践是维护一个计划池，对常用尺寸提前申请计划并驻留复用。临时工作区显存默认由 muFFT 内部管理，也可自行接管：

mufftHandle plan = NULL;
size_t workSize;
mufftCreate(&plan);
mufftSetAutoAllocation(plan, 0);                      // 关闭库自动分配
mufftMakePlan1d(plan, Nx, MUFFT_C2C, 1, &workSize);    // 查询所需工作区大小
mufftSetWorkArea(plan, buffer + dataOffset - workSize); // 绑定用户自管理显存

单机多卡场景下多卡接口声明在 mufftXt.h：需用 mufftXtSetGPUs() 设置参与变换的 GPU；资源申请/释放改用 mufftXtMalloc()/mufftXtFree()；数据传输改用 mufftXtMemcpy()；执行变换调用 mufftXtExecDescriptorC2C()。

4.4 muSPARSE：四种稀疏存储格式

muSPARSE 针对大部分元素为 0 的稀疏矩阵做深度优化，支持稀疏-稠密、稀疏-稀疏多类运算及格式转换（如 bsr2csr）。以 4×4 矩阵 [[1,0,6,2],[1,6,0,0],[0,0,3,0],[0,0,3,4]] 为例，四种常见格式：

核心计算接口遵循统一模式，以 SpMV（稀疏矩阵-向量乘法，计算 y := alpha·op(A)·x + beta·y）为例：

musparseStatus_t musparseSpMV(
    musparseHandle_t handle, musparseOperation_t op, const void* alpha,
    const musparseSpMatDescr_t matA, const musparseDnVecDescr_t vecX,
    const void* beta, const musparseDnVecDescr_t vecY,
    musaDataType computeType, musaSpMVAlg_t alg,
    size_t* buffer_size, void* externalBuffer);

matA 由 musparseCreateCsr() 创建，vecX/vecY 由 musparseCreateDnVec() 创建；先调用一次获取 buffer_size，再手动分配 externalBuffer。SpGEMM（稀疏矩阵-矩阵乘法）在此基础上多了 matB/matD/matC 三个稀疏矩阵描述符与 stage 参数（表示算法执行到"确定输出结构"还是"计算数值"阶段）；SpSV 求解稀疏三角方程组；SpVV 计算两个稀疏向量的内积。这五类接口都遵循"创建描述符 → 查询/分配外部缓冲区 → 调用计算接口"的统一模式。muSPARSE 还提供无主元不完整 LU 分解 musparseXcsrilu0，使用前需先用对应的 _bufferSize() 接口获取临时内存大小，再用 _analysis() 分析矩阵是否符合计算要求。

4.5 muSOLVER：LU 分解求解线性方程组

muSOLVER 专注矩阵分解、线性方程组求解和特征值计算，支持 FP16/FP32/TF32/FP64 混合计算。稠密矩阵分解涵盖 LU、QR、Cholesky、SVD、特征值分解；稀疏矩阵分解（迭代法 GMRES、直接法 LU/Cholesky）据原书标注仍在开发中，尚未发布。基本流程：创建 handle → 创建 stream 并绑定 → 数据 Host→Device → 调用功能函数 → 结果 Device→Host → 释放资源。

#include 
#include 
#include 
#include 
#include 

int main(int argc, char *argv[])
{
    musolverHandle_t musolverH = NULL;
    musaStream_t stream = NULL;

    const int m = 3;
    const int lda = m;   // A 的主维跨度
    const int ldb = m;   // B 的主维跨度

    const std::vector A = {1.0,4.0,7.0, 2.0,5.0,8.0, 3.0,6.0,10.0}; // 列优先存储
    const std::vector B = {1.0, 2.0, 3.0};

    std::vector X(m, 0);
    std::vector LU(lda * m, 0);
    std::vector Ipiv(m, 0);
    int info = 0;

    float *d_A = nullptr, *d_B = nullptr, *d_work = nullptr;
    int *d_Ipiv = nullptr, *d_info = nullptr;
    int lwork = 0;
    const int pivot_on = 0;  // 是否启用选主元

    // 步骤 1：创建句柄并绑定非阻塞流
    musolverCreate(&musolverH);
    musaStreamCreateWithFlags(&stream, musaStreamNonBlocking);
    musolverSetStream(musolverH, stream);

    // 步骤 2：分配显存并异步拷贝
    musaMalloc(reinterpret_cast(&d_A), sizeof(float) * A.size());
    musaMalloc(reinterpret_cast(&d_B), sizeof(float) * B.size());
    musaMalloc(reinterpret_cast(&d_Ipiv), sizeof(int) * Ipiv.size());
    musaMalloc(reinterpret_cast(&d_info), sizeof(int));
    musaMemcpyAsync(d_A, A.data(), sizeof(float) * A.size(), musaMemcpyHostToDevice, stream);
    musaMemcpyAsync(d_B, B.data(), sizeof(float) * B.size(), musaMemcpyHostToDevice, stream);

    // 步骤 3：查询 LU 分解所需工作区大小
    musolverSgetrf_bufferSize(musolverH, m, m, d_A, lda, &lwork);
    musaMalloc(reinterpret_cast(&d_work), sizeof(float) * lwork);

    // 步骤 4：执行 LU 分解
    if (pivot_on) {
        musolverSgetrf(musolverH, m, m, d_A, lda, d_work, d_Ipiv, d_info);
    } else {
        musolverSgetrf(musolverH, m, m, d_A, lda, d_work, NULL, d_info); // 不选主元传 NULL
    }

    if (pivot_on) {
        musaMemcpyAsync(Ipiv.data(), d_Ipiv, sizeof(int) * Ipiv.size(), musaMemcpyDeviceToHost, stream);
    }
    musaMemcpyAsync(LU.data(), d_A, sizeof(float) * A.size(), musaMemcpyDeviceToHost, stream);
    musaMemcpyAsync(&info, d_info, sizeof(int), musaMemcpyDeviceToHost, stream);
    musaStreamSynchronize(stream);

    if (0 > info) { printf("%d-th parameter is wrong \n", -info); exit(1); }

    // 步骤 5：求解 A * X = B
    if (pivot_on) {
        musolverSgetrs(musolverH, MUBLAS_OP_N, m, 1, d_A, lda, d_Ipiv, d_B, ldb, d_info);
    } else {
        musolverSgetrs(musolverH, MUBLAS_OP_N, m, 1, d_A, lda, NULL, d_B, ldb, d_info);
    }
    musaMemcpyAsync(X.data(), d_B, sizeof(float) * X.size(), musaMemcpyDeviceToHost, stream);
    musaStreamSynchronize(stream);

    // 步骤 6：释放资源
    musaFree(d_A); musaFree(d_B); musaFree(d_Ipiv); musaFree(d_info); musaFree(d_work);
    musolverDestroy(musolverH);
    musaStreamDestroy(stream);
    musaDeviceReset();
    return 0;
}

预期结果：B = [1, 2, 3]^T，解得 X ≈ [-0.3333, 0.6667, 0]^T。该示例体现"分解（getrf）→ 求解（getrs）"两阶段范式，以及"查询工作区大小 → 分配 → 计算"这一在 MUSA 高性能库中反复出现的资源管理套路。

五、AI 与通信库：muDNN 与 MCCL

5.1 muDNN：张量与算子对象模型

muDNN 是面向深度神经网络的 GPU 加速原语库，提供 C++ 风格接口，覆盖注意力、卷积、矩阵乘法、归一化、Softmax、池化、TopK、归约等运算。库文件 libmudnn.so 由三部分构成：mudnn_xmma（注意力/卷积/矩阵乘法）、mudnn_tensor（张量基础运算）、mudnn_ops（Batch Normalization、Softmax、Dropout 等常用算子）。

Handle 代表一个计算上下文环境，管理设备号与 MUSA 计算流，并可控制 TF32 数据类型的使用：

#include "mudnn.h"
#include "musa_runtime.h"

// 创建 muDNN Handle，默认使用 0 号设备
::musa::dnn::Handle handle;
// 或指定设备号（使用 1 号设备）
// ::musa::dnn::Handle handle(1);

int device_id = handle.GetDeviceId();

musaStream_t stream1;
musaStreamCreate(&stream1);
handle.SetStream(stream1);

musaStream_t current_stream = handle.GetStream();

handle.SetAllowTF32(true);
bool allow_tf32 = handle.GetAllowTF32();

MemoryHandler（类型定义为 std::unique_ptr<void, std::function<void(void*)>>）负责管理特定内存块生命周期；MemoryMaintainer（类型定义为 std::function<MemoryHandler(size_t)>）负责动态分配内存并返回 MemoryHandler。两者关系是"分配与释放职责分离"：MemoryMaintainer 负责分配，MemoryHandler 负责释放，本质是 C++ 中"工厂函数返回 RAII 智能指针"模式在 GPU 内存管理场景的应用。

#include "mudnn.h"
#include "musa_runtime.h"

void MemFree(void* ptr)
{
    if (ptr) { musaFree(ptr); }
}

::musa::dnn::MemoryHandler MemoryFunc(size_t size)
{
    void* data = nullptr;
    if (size) { musaMalloc(&data, size); }
    return ::musa::dnn::MemoryHandler(data, MemFree);
}

::musa::dnn::Handle handle;
::musa::dnn::Softmax op;
::musa::dnn::Tensor out;
::musa::dnn::Tensor in;
int dim = 0; // 需按实际维度、算法与模式设置
op.SetDim(dim);
op.SetAlgorithm(::musa::dnn::Softmax::Algorithm::DIRECT);
op.SetMode(::musa::dnn::Softmax::Mode::SOFTMAX);
op.Run(handle, out, in, MemoryFunc);

muDNN 常见状态码：

5.2 MCCL：多卡通信原语

MCCL（Moore Threads Collective Communications Library）提供 GPU 间通信原语，用于多 GPU、多节点通信，支持拓扑感知能力。支持的通信原语：all-gather、all-reduce、broadcast、reduce、reduce-scatter、p2p send/receive（可组合实现 alltoall/gather/scatter）。

底层实现基于 MUSA 内存复制和核函数完成局部归约，可利用节点内 PCIe 和 MTLink 高速互联，以及节点间 RDMA 网络。节点内数据传输核心技术是 GPUDirect P2P（避免数据经 CPU 中转）；跨节点核心技术是 GPUDirect RDMA（数据无需经 CPU/系统内存中转）。MCCL 提供标准 C 语言 API，接口标准与 MPI 一致，但额外增加了 stream 参数以便与 MUSA 编程模型集成同步。

// 模式一：单进程/线程负责单设备
mcclCommInitRank(&comm, n_ranks, unique_id, my_rank);
mcclAllReduce(sendbuff, recvbuff, count, datatype, op, comm, stream);
musaStreamSynchronize(stream);

// 模式二：单进程/线程负责多设备，必须用 Group Calls 包裹，避免死锁
mcclGroupStart();
for (int i = 0; i < n_local_devs; i++) {
    musaSetDevice(devices[i]);
    mcclCommInitRank(comms + i, n_ranks, unique_id, ranks[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < n_local_devs; i++) {
    mcclAllReduce(sendbuff[i], recvbuff[i], count, datatype, op, comms[i], streams[i]);
}
mcclGroupEnd();

for (int i = 0; i < n_local_devs; i++) {
    musaStreamSynchronize(streams[i]);
}

六、性能分析工具（概览）

Moore Perf Tools 是一套低开销性能分析工具，通过全栈时间线可视化、低开销数据采集和跨硬件分析，帮助开发者定位 MUSA 及 Graphics 应用程序的性能瓶颈，适用于深度学习、HPC、游戏开发等领域。包含两个子工具：Moore Perf System（系统级、自上而下的性能问题定位）与 Moore Perf Compute（计算侧性能分析）。详细使用方法将在后续章节（对应 m11-l10）展开，本课仅做概览性提及。

七、Step-by-Step：三步动手实验

实验一：CUDA 代码到 MUSA 代码的等效转换

1. 阅读原始 CUDA 代码，识别 CUDA 专属 API：
   cuda_runtime.h、__global__ 内核声明、cudaMalloc、cudaMemcpy、
   <<>> 内核启动语法、cudaFree

2. 按 Runtime API 命名习惯替换：
   cuda_runtime.h  -> musa_runtime.h
   cudaMalloc      -> musaMalloc
   cudaMemcpy      -> musaMemcpy
   cudaFree        -> musaFree
   cudaMemcpyHostToDevice -> musaMemcpyHostToDevice（方向常量同理）

3. 内核函数声明（__global__）与内核启动语法（<<<>>>）在 MUSA 中通常保持一致或近似
   （具体是否完全一致以 MUSA 编程模型章节描述为准，
    〔存疑，需结合编程模型章节确认，不做臆测〕）

4. 用 mcc 编译转换后的 .mu 文件：
   mcc target.mu -lmusart -L/usr/local/musa/lib -O2 -o target

5. 运行可执行文件，核对每个元素误差：
   fabs(h_C[i] - 3.0f) > 1e-5 是否成立，全部通过则视为验证成功

6. 记录 API 差异点，形成"CUDA -> MUSA 对照表"，便于后续项目迁移

实验二：muBLAS 矩阵乘法性能对比

1. 用 muBLAS 矩阵乘法函数（单精度 Sgemm / 双精度 Dgemm）
   实现 M x K 矩阵 A 与 K x N 矩阵 B 相乘，得 M x N 矩阵 C

2. 设置三组规模：(128,128,128)、(1024,1024,1024)、(4096,4096,4096)
   记录 muBLAS 实现的运行时间与 GFLOPS

3. 自行编写 MUSA 核函数实现矩阵乘法（不借助 muBLAS），
   可从最朴素的三重循环开始，再尝试分块（tiling）与共享内存优化

4. 对比两者执行效率差异：通常规模越大，muBLAS 的寄存器分块、
   共享内存复用、指令级并行等深度优化优势越明显

5. 将三组规模的运行时间和 GFLOPS 整理成表格或图表，
   直观呈现"规模增大后库函数优势扩大"的趋势

实验三：基于 muDNN 的 MNIST 分类任务

1. 调用 muDNN 的张量、卷积层、全连接层、Softmax 等算子，
   构建"一个卷积层 + 一个池化层 + 两个全连接层"的分类模型
   （调用范式：创建算子对象 -> SetXxx() 配置参数 -> Run() 执行）

2. 加载 MNIST 数据集，完成归一化、reshape 等预处理，
   确保输入张量维度、数据类型与算子要求匹配（否则触发 INVALID_PARAMETER）

3. 实现前向传播、交叉熵损失计算与反向传播（调用梯度算子）

4. 训练模型 10 个 epoch，评估测试集准确率

5. 调整卷积核大小或全连接层神经元数量，观察性能变化

6. 全程用 MemoryHandler/MemoryMaintainer 机制管理张量显存，
   并在每次算子调用后检查返回状态码

八、常见陷阱与排错提示

• 列优先 vs 行优先存储混淆：muBLAS 采用列优先存储，直接套用行优先思维初始化矩阵或计算主维跨度会导致结果错误或访问越界。
• Driver API 与 Runtime API 上下文语义混用不当：不了解"Runtime API 隐式创建 Primary Context"规则，在多线程/多 GPU 场景容易上下文管理混乱。
• muFFT inplace 模式的显存对齐留白：忽略输入/输出共用显存时的对齐留白（padding），按"紧密排布"读取会导致数据错位。
• 稀疏矩阵格式选择不当：COO/CSR/CSC/BSR 适合不同访问模式，格式选错可能导致性能不升反降。
• muSOLVER 中忽视选主元对数值稳定性的影响：为性能关闭选主元可能在部分矩阵上产生数值不稳定的结果。
• MCCL 单线程多卡未使用 Group Calls：容易发生死锁，这是官方明确强调的必要条件。
• muDNN 状态码被忽略：只关注计算结果、不检查状态码，会让参数错误、架构不匹配等问题难以定位。

九、版本与兼容性说明

十、本章知识与后续课程

• m11-l2（第 2 章）：硬件架构——本课的 TCE/TME 由 muDNN、muBLAS 在软件层封装调用
• m11-l3（第 3 章）：编程模型——Kernel/Context/Stream 是本课 Driver/Runtime API 的具体编程接口
• m11-l9（第 7 章）：MUSA 程序调试——muDNN 状态码、编译期属性提示与调试信息在排错时会反复用到
• m11-l10（第 9 章）：性能分析工具——Moore Perf System / Moore Perf Compute 的详细使用方法