C++ 语言支持

MUSA C++ 语言扩展和特性

概述

MUSA SDK 通过 MCC 编译器提供了对 C++ 语言的全面支持。MCC 基于 LLVM 构建，继承了 LLVM 对 C++ 标准的完整支持，并添加了针对 MUSA 设备的特定扩展。当前支持的 LLVM 版本是 LLVM 14。

本附录详细介绍 MCC 编译器支持的 C++ 语言特性，包括语言扩展、内置函数和编程接口。

---

C++ 标准支持

标准版本支持

MCC 支持以下 C++ 标准：

标准	支持状态	说明
C++11	✓ 完全支持	包括 auto、lambda、decltype 等
C++14	✓ 完全支持	包括泛型 lambda、返回类型推导等
C++17	✓ 完全支持	包括结构化绑定、constexpr lambda 等
C++20	✓ 部分支持	概念（concepts）、协程（实验性）

核心语言特性

以下代码用于说明语言特性写法。实际编译时，请根据所选 C++ 标准补齐头文件、变量和函数定义。

// C++11 特性
auto ptr = make_unique<int>(42);
auto lambda = [](int x) { return x * 2; };
auto [a, b] = getPair();  // C++17 结构化绑定

// constexpr 和 consteval
constexpr int square(int x) { return x * x; }
consteval int fast_square(int x) { return x * x; }

// 概念 (C++20)
template<typename T>
concept Addable = requires(T a, T b) { a + b; };

template<Addable T>
T add(T a, T b) { return a + b; }

---

MUSA 特定扩展

设备函数限定符

MCC 提供设备端代码限定符：

限定符	描述
__device__	在设备上执行的函数
__host__	在主机上执行的函数（默认）
__global__	从主机调用、在设备上执行的函数（内核）
__device__ __host__	同时支持设备和主机的函数

// 设备函数 - 只在设备上执行
__device__ float device_func(float x) {
    return x * x;
}

// 全局函数 - 从主机调用
__global__ void kernel_func(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = device_func(data[idx]);
    }
}

// 主机函数 - 只在主机执行
__host__ void host_only_func() {
    // 纯主机代码
}

// 设备/主机兼容函数
__device__ __host__ float compatible_func(float x) {
    return sin(x);  // MUSA 提供 sin 的设备版本
}

内置函数

MCC 提供丰富的设备内置函数：

数学函数

// 基础数学函数
__device__ float sin(float x);
__device__ float cos(float x);
__device__ float tan(float x);
__device__ float exp(float x);
__device__ float log(float x);
__device__ float sqrt(float x);
__device__ float rsqrt(float x);  // 快速平方根倒数
__device__ float pow(float x, float y);

// 快速数学函数 (低精度高性能)
__device__ float __sinf(float x);   // 快速正弦
__device__ float __expf(float x);   // 快速指数
__device__ float __powf(float x, float y);  // 快速幂

内存操作

// 加载/存储函数
template<typename T>
__device__ T __ldg(const T *ptr);  // 只读缓存加载

template<typename T>
__device__ void __st(void *ptr, T value);  // 释放存储

// 异步内存操作
__device__ void __threadfence();  // 线程同步
__device__ void __threadfence_block();  // 块内同步

整数和位操作

// 整数内置函数
__device__ int __mul24(int x, int y);  // 24位整数乘法
__device__ long long __mul64hi(long long x, long long y);  // 64位乘法高位
__device__ int __clz(int x);  // 前导零计数
__device__ int __popc(int x);  // 人口计数（1位数）
__device__ int __ffs(int x);  // 首个置位位置

内置变量

MCC 提供预定义的内置变量：

// 线程和块索引
blockIdx.x, blockIdx.y, blockIdx.z  // 块索引
threadIdx.x, threadIdx.y, threadIdx.z  // 线程索引
blockDim.x, blockDim.y, blockDim.z  // 块维度
gridDim.x, gridDim.y, gridDim.z  // 网格维度

---

模板支持

设备端模板

// 模板设备函数
template<typename T>
__device__ T add(T a, T b) {
    return a + b;
}

// 模板全局函数
template<typename T>
__global__ void vector_add(T *a, T *b, T *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = add(a[idx], b[idx]);
    }
}

// 实例化
vector_add<float><<<grid, block>>>(d_a, d_b, d_c, n);
vector_add<double><<<grid, block>>>(d_a, d_b, d_c, n);

模板参数推导

// 自动模板参数推导
template<typename T>
__global__ void kernel(T *data, int n) {
    // 使用 T
}

kernel(d_data, n);  // 自动推导 T

---

标准库支持

Thrust 库

MCC 支持 Thrust 并行算法库：

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>

// 设备向量
thrust::device_vector<float> d_vec(n);
thrust::host_vector<float> h_vec(n);

// 并行排序
thrust::sort(d_vec.begin(), d_vec.end());

// 变换
thrust::transform(d_vec.begin(), d_vec.end(), 
    d_vec.begin(), thrust::negate<float>());

标准容器

// 设备端可以使用部分 STL 容器
#include <vector>

__global__ void process_vector(const float *data, int n) {
    // 在设备代码中直接访问指针
}

---

语言特性兼容性

Lambda 表达式

// 设备端 lambda (C++14+)
auto lambda = [] __device__ (int x) { return x * 2; };

// 如需在设备端使用外部值，建议通过 kernel 参数或普通函数参数传入，
// 不要捕获 host 变量。

---

性能优化特性

谓词执行

// 使用谓词避免分支
__global__ void pred_kernel(float *data, float *result, float threshold) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 编译器自动使用谓词
    if (data[idx] > threshold) {
        result[idx] = data[idx] * 2.0f;
    } else {
        result[idx] = data[idx];
    }
}

---

编译选项

语言标准选择

# 指定 C++ 标准
mcc -std=c++17 -o app app.cpp

# 启用 C++20 语言模式
mcc -std=c++20 -o app app.cpp

设备代码编译

# 编译设备代码
mcc --musa-device-only -c -o device.o device.mu

# 完整编译
mcc --offload-arch=<arch> app.mu -lmusart -L/usr/local/musa/lib -o app

---

与 CUDA C++ 的差异

MCC 尽力保持与 nvcc 的兼容性，但存在一些差异：

特性	nvcc	MCC
语言标准	C++17 默认	C++14 默认
运行时	CUDA Runtime	MUSA Runtime
头文件	cuda_runtime.h	musa_runtime.h

---

相关文档

• MUSA SDK 概述
• MUSA C++ 扩展
• MCC 编译器用户手册
• 编程指南概述