MUSA C++ 语言扩展
概述
MUSA C++ 扩展是 MCC 编译器为 MTGPU 编程提供的语言扩展。它保持了与标准 C++ 的兼容性，同时添加了用于表达并行计算的语法结构。
核心扩展内容：
- 函数限定符：
__global__,__device__,__host__ - 内存限定符：
__shared__,__constant__,__managed__ - 内置变量：
threadIdx,blockIdx,gridDim,blockDim - Kernel 启动语法：
<<<grid, block>>>
本文档详细介绍 MUSA C++ 扩展的语法和用法。
变量限定符
device 限定符
__device__ 限定符声明的函数在 MUSA 设备上执行：
__device__ float device_function(float x) {
return x * x + 1.0f;
}
host 限定符
__host__ 限定符声明的函数在 MUSA 主机上执行：
__host__ float host_function(float x) {
return std::sqrt(x);
}
global 限定符
__global__ 限定符声明从主机调用、在设备上执行的函数（核函数）：
__global__ void vector_add_kernel(
const float *a,
const float *b,
float *c,
int n
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
device host 组合
同时使用两个限定符可以使函数同时支持设备和主机：
__device__ __host__ float compatible_function(float x) {
return x * 2.0f;
}
内置变量
MCC 提供了预定义的内置变量：
线程索引
// 线程在块内的索引
threadIdx.x // 0 到 blockDim.x - 1
threadIdx.y // 0 到 blockDim.y - 1
threadIdx.z // 0 到 blockDim.z - 1
块索引
// 块在网格中的索引
blockIdx.x // 0 到 gridDim.x - 1
blockIdx.y // 0 到 gridDim.y - 1
blockIdx.z // 0 到 gridDim.z - 1
块维度
// 块的维度（线程数）
blockDim.x
blockDim.y
blockDim.z
网格维度
// 网格的维度（块数）
gridDim.x
gridDim.y
gridDim.z
Warp 信息
// 使用设备提供的实际 Warp 大小，不要硬编码产品或固定值
warpSize
内置函数概览
MUSA 提供丰富的设备端内置函数，详细说明请参阅 原子函数 和 Warp 函数。
同步函数
| 函数 | 作用域 | 说明 |
|---|---|---|
__syncthreads() | 线程块 | 块内所有线程的同步屏障 |
__threadfence() | 设备级 | 约束设备范围内的内存访问顺序 |
__threadfence_block() | 线程块 | 确保块内所有内存写操作可见 |
__threadfence_system() | 系统级 | 约束系统范围内的内存访问顺序 |
__syncwarp() | 线程束 | 线程束级同步屏障 |
当不同 block 用非 atomic 写相同的 global memory 地址, 需要在写回语句后加一个 __threadfence() 调用来确保正确性(目标架构小于等于 mp_31), 示例如下:
__global__ void func(int* dst){
...
*dst = block_id;
__threadfence();
...
}
原子函数
| 函数 | 说明 |
|---|---|
atomicAdd(), atomicSub() | 原子加减 |
atomicExch() | 原子交换 |
atomicInc(), atomicDec() | 原子递增递减 |
atomicAnd(), atomicOr(), atomicXor() | 原子位运算 |
atomicCAS() | 原子比较交换 |
原子函数支持 _block 和 _system 变体，分别用于块级和系统级作用域。
Warp 级函数
| 函数 | 说明 |
|---|---|
__shfl_sync() | Warp 内线程间数据交换 |
__all_sync(), __any_sync() | Warp 投票 |
__ballot_sync() | 位选投票 |
__reduce_*_sync() | Warp 归约操作 |
数学函数
// 标准精度
sin(x), cos(x), exp(x), log(x), sqrt(x)
// 快速版本（低精度，高性能）
__sinf(x), __cosf(x), __expf(x), __logf(x), __sqrtf(x)
| 函数类型 | 精度 | 性能 | 适用场景 |
|---|---|---|---|
| 标准函数 | 高 | 中 | 精确计算 |
| 快速函数 | 低 | 高 | 性能敏感场景 |
核函数配置
执行配置
核函数启动使用 <<<grid, block, shared_mem, stream>>> 语法：
// grid: 网格维度（块数）
// block: 块维度（每块线程数）
// shared_mem: 动态共享内存大小（字节）
// stream: 执行流
kernel_name<<<grid_dim, block_dim, shared_mem, stream>>>(args);
示例
// 一维配置
int n = 1024;
int block_size = 256;
int grid_size = (n + block_size - 1) / block_size;
kernel<<<grid_size, block_size>>>(args);
// 二维配置（图像处理）
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (height + 15) / 16);
kernel<<<grid, block>>>(args);
内存模型
MUSA 支持多种内存类型。本节介绍设备端代码（Device Code）中常用的存储空间和访问方式。
全局内存
// 在全局内存中声明变量
__device__ float global_array[1024];
// 动态分配全局内存（Runtime API）
float* d_ptr;
musaMalloc(&d_ptr, size);
共享内存
// 静态共享内存
__global__ void kernel1() {
__shared__ float shared_mem[256];
// 使用 shared_mem
}
// 动态共享内存（通过核函数配置传递）
__global__ void kernel2() {
extern __shared__ float shared_mem[];
// 使用 shared_mem
}
// 启动时指定大小
kernel2<<<grid, block, 256 * sizeof(float)>>>(args);
常量内存
// 声明常量内存
__constant__ float const_array[1024];
// 从主机初始化
musaMemcpyToSymbol(const_array, h_data, size);
纹理内存和表面函数的创建、资源描述和生命周期由 Runtime API 管理。
编译
编译 MUSA C++ 程序时，显式指定已发布目标架构。当前示例使用 mp_31：
mcc -mtgpu --offload-arch=mp_31 vector_add.mu -o vector_add
.mu 是 MUSA C++ 源文件后缀。使用模板、lambda 或 C++ 标准库组件时，可显式指定 C++ 标准：
mcc -mtgpu --offload-arch=mp_31 -std=c++14 kernel.mu -o app
编译或性能比较前，固定 SDK、驱动、编译器、目标设备、C++ 标准、优化选项和数值相关选项。对不满足目标架构、数据类型或形状条件的路径，保留普通 Device Code 回退路径，并以公开头文件和编译测试确认实际支持情况。
编译指示
循环展开
#pragma unroll
for (int i = 0; i < 8; i++) {
// 编译器展开循环， 提高性能
}
强制内联
__forceinline__ float fast_add(float a, float b) {
return a + b;
}
错误处理
MUSA 使用 musaError_t 类型表示错误。错误枚举和数值应以 SDK 的 Runtime API 头文件为准，不要在应用文档中重声明或复制不完整的枚举。
错误处理示例：
musaError_t err = musaMalloc(&ptr, size);
if (err != musaSuccess) {
fprintf(stderr, "Memory allocation failed: %d\n", err);
return -1;
}
扩展关键字总览
| 关键字 | 描述 |
|---|---|
__device__ | 设备函数 |
__host__ | 主机函数 |
__global__ | 核函数 |
__shared__ | 共享内存 |
__constant__ | 常量内存 |
__managed__ | 统一内存 |
__restrict__ | 指针非别名 |
__inline__ | 内联函数 |
__noinline__ | 不内联 |
