第一个MUSA程序 | 摩尔线程 GPU 学习中心

一、概述

在本课中，我们将编写并运行你的第一个完整的MUSA GPU程序——向量加法（Vector Addition）。这是GPU编程的"Hello World"，虽然简单，但它涵盖了GPU编程的所有核心概念：内存分配、数据传输、Kernel编写和启动。

通过这个程序，你将学会：

如何编写GPU Kernel函数（用 __global__ 标记）
如何在GPU上分配和释放内存（musaMalloc / musaFree）
如何在CPU和GPU之间传输数据（musaMemcpy）
如何配置Grid和Block来启动Kernel
如何编译和运行MUSA程序

二、异构计算基础

在GPU编程中，我们使用异构计算（Heterogeneous Computing）模型：程序的一部分在CPU上运行（称为Host端），另一部分在GPU上运行（称为Device端）。两者通过PCIe总线连接，各自拥有独立的内存空间。

2.1 Host与Device

Host（主机）：CPU及其内存。负责程序的总体控制流、文件I/O、网络通信等
Device（设备）：GPU及其显存。负责大规模并行计算任务

一个典型的GPU程序的执行流程：

在Host端准备数据（CPU内存）
在Device端分配显存（GPU显存）
将数据从Host复制到Device
在Device上启动Kernel进行并行计算
将计算结果从Device复制回Host
释放Device端显存

💡

为什么需要数据拷贝？

CPU和GPU有各自独立的内存空间（CPU用DDR内存，GPU用GDDR显存），它们不能直接访问对方的内存。因此数据需要通过PCIe总线在两者之间显式复制。MUSA也支持统一内存（Unified Memory），可以简化这一过程，但底层仍然涉及数据传输。

三、向量加法完整代码

下面是完整的向量加法MUSA程序。我们将逐行解读每一部分：

cpp

#include <musa_runtime.h>
#include <stdio.h>

// GPU Kernel函数：向量加法
// __global__ 表示此函数在GPU上执行，从CPU端调用
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    // 计算当前线程的全局索引
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    // 边界检查：确保不越界访问
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1024;                    // 向量长度
    size_t size = n * sizeof(float); // 数据大小（字节）

    // ========== 第1步：在Host端分配内存并初始化数据 ==========
    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_c = (float *)malloc(size);

    // 初始化输入数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }

    // ========== 第2步：在Device端分配显存 ==========
    float *d_a, *d_b, *d_c;
    musaMalloc(&d_a, size);
    musaMalloc(&d_b, size);
    musaMalloc(&d_c, size);

    // ========== 第3步：将数据从Host复制到Device ==========
    musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice);
    musaMemcpy(d_b, h_b, size, musaMemcpyHostToDevice);

    // ========== 第4步：配置Grid/Block并启动Kernel ==========
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // ========== 第5步：将结果从Device复制回Host ==========
    musaMemcpy(h_c, d_c, size, musaMemcpyDeviceToHost);

    // ========== 第6步：验证结果 ==========
    printf("Result: c[0]=%f, c[1]=%f, c[2]=%f\n", h_c[0], h_c[1], h_c[2]);
    // 预期输出: c[0]=0.000000, c[1]=3.000000, c[2]=6.000000

    // 验证所有结果
    int errors = 0;
    for (int i = 0; i < n; i++) {
        float expected = i * 1.0f + i * 2.0f;
        if (h_c[i] != expected) {
            errors++;
            if (errors <= 5) {
                printf("Error at index %d: got %f, expected %f\n",
                       i, h_c[i], expected);
            }
        }
    }
    if (errors == 0) {
        printf("✓ All %d elements correct!\n", n);
    } else {
        printf("✗ %d errors found.\n", errors);
    }

    // ========== 第7步：释放内存 ==========
    musaFree(d_a);
    musaFree(d_b);
    musaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

四、逐行代码解读

4.1 头文件引入

cpp

#include <musa_runtime.h>  // MUSA运行时API
#include <stdio.h>         // C标准输入输出

musa_runtime.h 是MUSA编程的核心头文件，提供了所有GPU管理和数据传输的API函数，如 musaMalloc、musaMemcpy、musaFree 等。

4.2 Kernel函数

cpp

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

关键点解析：

__global__：函数执行空间说明符，表示此函数在GPU上执行，从CPU端通过 <<<...>>> 语法调用
blockDim.x：当前Block中的线程数（本例中为256）
blockIdx.x：当前线程所在的Block索引
threadIdx.x：当前线程在Block内的索引
全局索引公式：i = blockDim.x * blockIdx.x + threadIdx.x 计算每个线程的唯一全局ID
边界检查：if (i < n) 防止最后一个Block中多余的线程越界访问

4.3 Grid和Block配置

cpp

int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

为什么选择256个线程？这是一个经验值。每个Block的线程数通常选择32的倍数（因为Warp大小为32），128-512之间通常能获得较好的性能。

Block数量计算：(n + threadsPerBlock - 1) / threadsPerBlock 是向上取整除法，确保有足够的线程覆盖所有元素。例如n=1024，threadsPerBlock=256，则blocksPerGrid=4。

Grid: blocksPerGrid = 4 个Block

每个Block: threadsPerBlock = 256 个Thread

总线程数: 4 × 256 = 1024 个Thread
每个Thread处理一个向量元素

图3：本例中的Grid/Block/Thread配置 — 4个Block × 256个Thread = 1024个并行线程

4.4 内存管理API详解

MUSA提供的内存管理API与CUDA非常相似：

cpp

// 在GPU上分配显存
// 类似于CPU上的malloc，但分配的是GPU显存
musaMalloc(&d_a, size);  // d_a指向GPU显存

// 数据传输：从CPU内存复制到GPU显存
musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice);
//         目标  源   大小  方向（Host→Device）

// 数据传输：从GPU显存复制回CPU内存
musaMemcpy(h_c, d_c, size, musaMemcpyDeviceToHost);
//         目标  源   大小  方向（Device→Host）

// 释放GPU显存
musaFree(d_a);  // 类似于CPU上的free

⚠️

内存管理注意事项

1. 不要在CPU端解引用GPU指针：d_a是GPU显存指针，在CPU代码中不能直接读写（如 *d_a = 5; 会导致段错误）。
2. 务必检查错误返回值：生产代码应检查每个MUSA API调用的返回值。
3. 记得释放显存：GPU显存是有限资源，用完必须释放。

五、编译与运行

5.1 编译

MUSA程序使用 mcc 编译器编译，使用方法类似于gcc/nvcc：

bash

# 保存代码为 vector_add.mu
# .mu 是MUSA源文件的扩展名（类似CUDA的.cu）

# 编译
mcc vector_add.mu -o vector_add

# 带优化编译
mcc -O2 vector_add.mu -o vector_add

# 编译时指定GPU架构（可选）
mcc --gpu-arch=mp_21 vector_add.mu -o vector_add

5.2 运行

bash

# 运行程序
./vector_add

# 预期输出：
# Result: c[0]=0.000000, c[1]=3.000000, c[2]=6.000000
# ✓ All 1024 elements correct!

5.3 调试技巧

bash

# 检查MUSA API错误的辅助宏（建议加入所有项目）
# 在代码开头添加：

cpp

#define MUSA_CHECK(call) \
do { \
    musaError_t err = call; \
    if (err != musaSuccess) { \
        fprintf(stderr, "MUSA error at %s:%d: %s\n", \
                __FILE__, __LINE__, musaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

// 使用示例：
MUSA_CHECK(musaMalloc(&d_a, size));
MUSA_CHECK(musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice));

六、练习与扩展

完成了向量加法程序后，尝试以下练习来加深理解：

修改向量大小：将n改为1,000,000（一百万），观察程序是否仍然正确运行。计算此时需要多少个Block？
实现向量乘法：修改Kernel函数，将 c[i] = a[i] + b[i] 改为 c[i] = a[i] * b[i]。
实现向量缩放：编写一个新Kernel vectorScale(float *a, float *out, float scale, int n)，将每个元素乘以一个标量。
添加计时：使用 musaEvent_t 来测量Kernel的执行时间。
尝试不同Block大小：将threadsPerBlock分别设置为32、64、128、256、512，观察对性能的影响。

✅

本课小结

本课通过一个完整的向量加法程序，学习了MUSA GPU编程的基本流程：Host端准备数据 → Device端分配显存 → 数据传输 → Kernel启动 → 结果回传 → 资源释放。掌握了__global__函数、musaMalloc/musaMemcpy/musaFree等核心API，以及Grid/Block/Thread的线程组织方式。这些是所有MUSA程序的基础模式。