第3章精讲：MUSA 编程模型（Kernel / Context / Stream / 内存管理）

一、内核函数：设备端并行任务的最小定义单元

MUSA 内核函数是在设备端（GPU）上执行的并行计算函数，是 MUSA 编程模型的核心构造。基本定义语法如下：

__global__ void kernelProcess(args...)
{
    // 内核实现
}

内核函数有三个关键特性需要牢记：第一，函数修饰符 __global__ 表示该函数在设备端上执行，可以从主机端调用，且返回类型必须为 void；第二，执行模型上遵循 SIMT 模式，大量线程同时执行相同代码，通过线程 ID 索引值区分每个线程身份及其需要处理的数据；第三，调用语法使用三重尖括号，这是 MUSA/CUDA 系语言的标志性语法：

kernelProcess<<>>(args...);

当用户启动一个 Kernel 时，设备上会启动 gridDim.x × gridDim.y × gridDim.z 个线程块，每个线程块包含 blockDim.x × blockDim.y × blockDim.z 个线程，每个线程根据自身索引处理不同区域的数据。下面是一份完整的向量加法示例，是本章最基础也最典型的入门代码：

__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    // 计算全局线程索引
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N)
    {
        C[i] = A[i] + B[i];
    }
}

// 调用：blockDim设置为1024（一个Block的最大线程数），gridDim设置为上取整(N/1024)
int threadsPerBlock = 1024;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<>>(d_A, d_B, d_C, N);

二、设备、上下文与任务流：三个容易混淆的概念

1. MUSA 设备（Device）

在 CPU-GPU 异构编程模型中，MUSA 设备指摩尔线程 GPU 硬件，是 MUSA 编程模型中的计算核心。GPU 由多个 MP 组成，每个 MP 包含多个 MUSA 核心，能够并行执行大量线程。每个 GPU 有唯一 ID，可通过 API 查询计算能力、内存大小等参数，下面是一份完整的设备查询示例：

#include 
#include 

int main()
{
    int deviceCount;
    musaGetDeviceCount(&deviceCount);
    std::cout << "发现 " << deviceCount << " 个MUSA设备" << std::endl;

    for (int i = 0; i < deviceCount; i++)
    {
        musaDeviceProp props;
        musaGetDeviceProperties(&props, i);
        std::cout << "设备 " << i << ": " << props.name << std::endl;
        std::cout << "  计算能力: " << props.major << "." << props.minor << std::endl;
        std::cout << "  全局内存: " << props.totalGlobalMem / (1024 * 1024) << " MB" << std::endl;
        std::cout << "  多处理器数量: " << props.multiProcessorCount << std::endl;
    }
    return 0;
}

2. 上下文（Context）

MUSA 上下文是一个软件环境，包含特定设备的状态信息，类似于 CPU 中的"进程"概念，每个上下文独立管理内存分配释放、模块加载、流调度与同步边界。所有 MUSA 操作（核函数调用、内存拷贝等）均需在上下文内执行，上下文销毁时会自动释放关联的硬件资源。关于上下文与主机线程的关联，有几条规则需要记住：一个主机线程在某一时刻只能有一个当前生效的设备上下文；用户通过调用 muCtxCreate() 驱动 API 创建一个设备上下文，创建时它自动成为当前主机线程的活动上下文；若当前线程不处于合法上下文，绝大多数 MUSA API（不含设备枚举或上下文管理相关函数）将返回错误码 MUSA_ERROR_INVALID_CONTEXT；为支持同一上下文被多个用户模块共用，MUSA 提供"上下文使用量计数器"，不同模块加载使用同一上下文时须调用 muCtxAttach() 增加计数器，使用完毕调用 muCtxDetach() 减少计数器。

3. MUSA 任务流（Stream）

MUSA 任务流（简称"流"）是一系列按顺序执行的 MUSA 操作，这些操作在设备上锁步执行。当主机端启动一个 Kernel 时，默认情况下主机会等待 GPU 执行完成；使用流可以打破这种默认同步行为，让主机计算、设备计算和数据传输并行执行，不同流中的操作可以交错执行，主机代码可以在设备操作执行期间继续运行。

三、内存管理：MUSA 编程的性能核心

设备与流的管理解决了"任务如何调度"的问题，而并行任务的高效执行离不开合理的内存管理。CPU-GPU 异构内存模型可以简化为三个功能模块：MUSA Runtime 内存管理、GPU 内存子系统、驱动与 GPU 交互。典型异构程序执行流程分为四步：申请 GPU 可访问内存并写入预处理后的数据；CPU 在 MUSA 运行时启动 GPU 程序，GPU 执行期间完成内存 I/O 与计算；CPU 异步等待 GPU 程序执行完毕；最终把计算结果写入磁盘并归还分配资源。在保障程序正确性的前提下，性能优化的核心几乎等价于对这四个步骤的逐一优化。

MUSA 按照局部性原理对数据访存模式和内存布局进行硬件层面的优化：时间局部性指一份数据被引用后，短时间内很可能会被二次引用；空间局部性指一份数据所在的内存位置被引用，其附近的内存也可能被引用。聚焦 GPU 架构内部，MUSA 编程模型根据数据的可见性与性能需求，划分出多种功能属性明确的存储空间：

这种分层划分允许开发者根据数据局部性特征，精细控制数据在不同层级间的流动，是后文矩阵乘法性能优化的直接理论基础——从寄存器到全局内存，延迟递增、容量递增，性能优化的一大主线正是尽可能把高频复用的数据留在低延迟层级。

四、任务与同步：保证并行执行的正确性

1. 两大任务类型

MUSA 中任务类型主要分为内核函数调用与内存操作两大类。内核函数进一步分为全局函数（内核函数本身，从主机端启动、在设备端运行）与设备函数（只能从内核或其他设备函数调用）：

// 全局函数（内核函数）：从主机端启动，在设备端运行
__global__ void globalFunc()
{
}

// 设备函数：只能从内核或其他设备函数调用
__device__ void deviceFunc()
{
}

2. 同步机制：主机-设备端同步与设备端同步

在多线程或多 GPU 编程中，同步机制是保障数据一致性和执行顺序的关键。MUSA 提供的同步机制分为两大类：

主机-设备端同步包含三种手段。流同步通过 musaStreamSynchronize() 函数同步一个或多个流，对重叠计算和数据传输非常有用：

musaStream_t stream;
musaStreamCreate(&stream);
kernelFunction<<>>();
musaStreamSynchronize(stream);
musaStreamDestroy(stream);

事件同步通过 musaEvent_t 对象记录事件的发生，并在之后的代码中等待事件完成，对测量执行时间或确保事件间的发生顺序非常有用：

musaEvent_t start, stop;
musaEventCreate(&start);
musaEventCreate(&stop);
musaEventRecord(start);
kernelFunction<<>>();
musaEventRecord(stop);
musaEventSynchronize(stop);

float milliseconds = 0;
musaEventElapsedTime(&milliseconds, start, stop);

设备同步使用 musaDeviceSynchronize() 函数，能在一个设备上实现所有任务的同步，等待该设备上此前发起的全部任务完成。

设备端同步包含线程内同步、线程块间同步与协同组同步三种。线程内同步通过 __syncthreads() 函数在同一个线程块内的线程之间进行同步，该函数会阻塞调用它的线程，直到同一线程块中的所有线程都执行到该点：

__global__ void kernelFunction()
{
    // ... 部分计算，写入共享内存 ...
    __syncthreads();
    // ... 读取共享内存中其他线程写入的数据 ...
}

如果要跨多个线程块同步，可以使用原子操作实现，具体方式因算法而异；协同组同步（cooperative_groups）与 __syncthreads() 相比，允许仅同步部分线程（如一个线程束），而不是整个 Block，粒度更细。

五、综合案例：矩阵乘法的并行化与内存访问优化

矩阵乘法是本章的综合应用案例，具备四个特征使它成为学习 MUSA 编程模型的最佳载体：计算任务高度可并行，输出矩阵中每个元素的计算彼此独立；计算与数据访问模式清晰，每个输出元素都对应输入矩阵的一行与一列；对内存层次结构高度敏感，不同存储方式对性能影响显著；能够完整覆盖 GPU 编程模型的核心要素。两个矩阵相乘的数学表达为 C（M×N）= A（M×K）× B（K×N），矩阵 C 中每个元素的计算公式为 c[i][j] 等于 A 第 i 行与 B 第 j 列对应元素乘积之和，累加 K 次。由于矩阵 C 中不同元素的计算在逻辑上彼此独立、不存在数据依赖关系，这正是并行计算的基础：在 MUSA 编程模型中，把输出矩阵中的元素映射为并行执行的线程，所有线程执行同一内核函数代码，每个线程通过自身索引确定所负责的输出位置。

性能的关键不在计算公式本身，而在数据如何被访问：同一行的 A 元素会被多个输出元素反复使用，同一列的 B 元素也会被多个线程访问，若每次计算都直接从全局内存读取，将产生大量冗余访存。下面是一份朴素实现，直接从全局内存读取数据，没有做任何复用优化：

__global__ void matMulNaive(float* A, float* B, float* C, int N)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N)
    {
        float sum = 0.0f;
        for (int k = 0; k < N; k++)
        {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

在 MUSA GPU 模型中，矩阵乘法加速的实质是通过分层存储重构了数据的访存模式，核心优化体现在三个层面：第一，把不连续的访问转化为合并访存（Coalesced Access）——矩阵在显存中通常按行连续存储，若线程直接读取矩阵 B 的列元素会触发跨步长的零散访问，导致显存带宽利用率急剧下降，因此用一组线程将零散的行列数据以连续合并方式批量读入、转存至片上共享内存；第二，提高数据复用、缓解带宽瓶颈——数据进入共享内存后访问延迟接近寄存器量级，矩阵 A 的行与矩阵 B 的列可被同一线程块内多个线程反复调用，全局内存访问频次因此降低约 N 倍；第三，软件分块与硬件矩阵计算单元（教材中称为 TCE）的协同——软件侧通过共享存储将高频复用的数据子块准备就绪，硬件侧的专用矩阵计算单元在寄存器层级流水化吞吐算力。下面是共享内存分块（Tiled）版本的完整实现：

#define TILE_SIZE 16

__global__ void matMulTiled(float* A, float* B, float* C, int N)
{
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    int numTiles = (N + TILE_SIZE - 1) / TILE_SIZE;
    for (int t = 0; t < numTiles; t++)
    {
        int aCol = t * TILE_SIZE + threadIdx.x;
        int bRow = t * TILE_SIZE + threadIdx.y;

        // 边界判断：N不能被TILE_SIZE整除时避免越界读取
        tileA[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        tileB[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;

        __syncthreads();

        for (int k = 0; k < TILE_SIZE; k++)
        {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }

        __syncthreads();
    }

    if (row < N && col < N)
    {
        C[row * N + col] = sum;
    }
}

六、Step-by-Step：从设备查询到矩阵乘法加速比实测

本实验的目标是亲手实现并对比朴素全局内存版矩阵乘法与共享内存分块版矩阵乘法，验证"提高数据复用可以显著提升性能"这一核心结论。建议按以下顺序展开：

第 1 步：环境与设备查询——先运行第二节的设备管理示例代码，确认当前环境可用的 MUSA 设备数量、计算能力、全局内存大小与多处理器数量，为后续设置合理的 blockDim/gridDim 提供依据。

第 2 步：内核函数基础打底——先独立实现并跑通第一节的 VecAdd 向量加法内核，确认自己已经熟悉 <<>> 调用语法与全局线程索引的计算方法，这是后续二维矩阵索引计算的基础。

第 3 步：实现朴素矩阵乘法内核——使用本课第五节给出的 matMulNaive 完整代码，为一个 N×N 的方阵（建议先用 N=512 做验证，避免调试阶段等待过久）分配设备内存、拷贝数据、启动内核、拷回结果，并用简单的 CPU 端双重循环验证正确性。

int N = 512;
size_t bytes = N * N * sizeof(float);

float *d_A, *d_B, *d_C;
musaMalloc(&d_A, bytes);
musaMalloc(&d_B, bytes);
musaMalloc(&d_C, bytes);

musaMemcpy(d_A, h_A, bytes, musaMemcpyHostToDevice);
musaMemcpy(d_B, h_B, bytes, musaMemcpyHostToDevice);

dim3 blockDim(16, 16);
dim3 gridDim((N + blockDim.x - 1) / blockDim.x, (N + blockDim.y - 1) / blockDim.y);

matMulNaive<<>>(d_A, d_B, d_C, N);
musaDeviceSynchronize();

musaMemcpy(h_C, d_C, bytes, musaMemcpyDeviceToHost);

第 4 步：实现共享内存分块矩阵乘法内核——把 matMulTiled 替换 matMulNaive 重新编译运行，注意分块大小 TILE_SIZE 需要与 blockDim 保持一致（本课示例中均为 16×16）。

第 5 步：用事件同步分别计时两个版本——使用第四节的事件同步代码分别测量朴素版本与分块版本的执行时间，计算加速比：

musaEvent_t start, stop;
musaEventCreate(&start);
musaEventCreate(&stop);

musaEventRecord(start);
matMulTiled<<>>(d_A, d_B, d_C, N);
musaEventRecord(stop);
musaEventSynchronize(stop);

float ms = 0;
musaEventElapsedTime(&ms, start, stop);
printf("分块版本耗时: %f ms\n", ms);

第 6 步：验证存算比与访存受限的关系——按朴素版本的访存模式手动计算：每个输出元素需要读取 A 的一行（N 个 float）与 B 的一列（N 个 float），对应 2N 次乘加，存算比大致是一个不随 N 增长而降低的常数量级；而分块版本让每个从全局内存加载到共享内存的数据块被块内多个线程反复复用，全局内存访问频次降低约 N 倍，存算比显著下降。把这个理论推导与你在第 5 步实测得到的加速比做对比讨论，理解"访存受限（Memory-Bound）"与"计算受限（Compute-Bound）"的区别——这正是 GPU 峰值算力往往远高于全局内存带宽所能供给的数据速率时，为什么必须依赖共享内存复用才能真正发挥算力的根本原因。

七、版本与兼容性说明

八、本课小结