5. MUSA编程基础

使用MUSA编程可以让用户编写程序完成任意的计算任务，这些计算任务中的一些步骤适合多线程并发执行，在GPU中运行可以获得巨大的加速收益。MUSA让用户能方便地编写出在GPU上运行的程序，MUSA语言是C++语言的扩展，其中绝大多数语法和c++一致，在c++语言的基础上添加了少许的新特性，如新的函数说明符，新的内建数据类型，以及调用gpu端函数的新的三括号语法等。本章将首先介绍使用这些新特性编写完整MUSA程序的整体思路和范式， 接着详细介绍这些拓展的特性。编写好的MUSA代码，需要通过mcc编译器编译成为mtgpu能识别的机器代码才能在GPU上执行，本章最后给出mcc编译器的使用方法。

5.1. MUSA基础编程

MUSA软件环境允许开发者使用C++作为高级编程语言进行开发。通过MUSA编程模型、编译器及相关API，MUSA为熟悉C++语言的用户提供了轻松简便的方法来编写执行在GPU设备上的程序。MUSA C++由相关的C++语言扩展和运行时库组成，核心的C++语言扩展在编程模型中引入，该扩展允许程序员将kernel定义为C++函数，并在调用kernel时使用扩展的语法来指定grid和block的维度。任何包含这些扩展的源文件都需要使用MUSA编译器进行编译。运行时库提供在主机上执行的C和C++ 函数，用于分配和释放Device内存、在Host内存和Device内存之间传输数据、管理具有多个Device的系统等。

5.1.1. MUSA编程模型

跟其他同类型的编程模型一样，MUSA编程模型是异构模型。异构意味着不同的体系结构，将程序划分为适用于不同体系结构的多个部分并执行在相应体系结构的硬件上，可以有效地提高程序的执行效率和性能；但这也意味着不同体系结构相互之间不能识别对方的二进制代码，需要在代码中做显式的区分。为了方便异构模型的描述和使用，需要将不同的体系结构进行抽象。在MUSA编程模型中，CPU和GPU两者分别被抽象为Host和Device，执行相应的指令；其次，MUSA编程模型假设Host和Device在各自的DRAM中维护独立的内存空间，相应的内存分别被称为Host Memory和Device Memory。所有的MUSA程序都需要Host和Device的协同工作，为了实现这一目的，典型的MUSA程序中既包含执行在CPU上的Host代码，也包含执行在GPU上的Device代码，两种代码协同工作，完成程序要求的任务。Host和Device之间依靠MUSA Runtime通信，管理内存的分配和释放、数据拷贝以及执行流控制等。MUSA编程模型中，Host一般只有一个，Device可以有一个或多个。典型的MUSA程序执行流程如下：

1. 分配Host Memory，并进行数据初始化；
2. 分配Device Memory，并将数据从Host Memory拷贝到Device Memory上；
3. 调用MUSA的kernel函数，在Device上执行指定的任务；
4. 将Device Memory中的结果拷贝回Host Memory；
5. 释放Device和Host上分配的Memory。

5.1.2. kernel

以上流程中最重要的部分是调用MUSA的kernel函数来执行指定的任务，MUSA C++通过允许程序员定义被称为kernel的C++函数作为C++语法的扩展。kernel函数是MUSA中的重要概念，指的是在Host端调用、在Device端执行的函数；当用户调用kernel时，Device上会启动N个不同的MUSA线程并行执行N次，而不是跟普通的C++函数那样只执行一次。跟普通的C++函数相比，kernel函数需要用__global__符号声明，并且返回值类型必须为void；调用kernel时需要用<<

>>来指定kernel需要执行的block数量和线程数量划分，注意"<<<"和">>>"中间不能有空格。同一任务中的每个线程都执行相同的kernel函数，并且每个线程会被分配到唯一的线程ID，该ID值可以通过内置的变量来获得。kernel的调用方式如下所示。

kernel<<<gridSize, blockSize>>>(prams...);

作为说明，以下代码中会用到内置变量threadIdx，将两个大小为N的向量x和y相加，并将结果存储到向量z中。

// Kernel definition
__global__ void vecAdd(float* x, float* y, float* z) {
    int i = threadIdx.x;
    z[i] = x[i] + y[i];
}

int main() {
    ...
    // Kernel invocation with N threads
    vecAdd<<<1, N>>>(x, y, z);
    ...
}

在这个例子中，Device上的N个VecAdd()线程对向量x、y和z的每个元素都执行了一组加法，注意x、y和z三个指针都指向GPU内存中的数据。threadIdx的含义是线程在当前block中的索引。

5.1.3. 线程层次结构

kernel执行的block和线程数量划分是MUSA程序实现高效运行的关键之一，要理解这一点，需要对kernel的线程层次结构有清晰的认识。我们知道，GPU上可以并行执行大量轻量级的线程，这是由GPU的硬件特性决定的，用户需要对这些线程进行层次化的组织，以充分利用硬件特性。

kernel在Device上执行时会启动若干个线程，单个kernel启动的所有线程被称为grid，同一个grid上的所有线程共享相同的全局内存空间，这是线程结构的第一层。而grid又可以分为一个或者多个thread block，每个block里面包含一个或多个线程，这些线程共享同一个MP内的资源，这是第二层。线程两层组织结构如下图所示，图中的gird和block均为2维的线程组织。

对于图1中的结构(水平方向为x轴)，其中的grid和block大小如以下代码中的gridSize和blockSize变量所示，

dim3 gridSize(3, 2);
dim3 blockSize(5, 3);
kernel<<<gridSize, blockSize>>>(prams...);

其中，gridSize和blockSize分别表示该grid内block的数量划分和单个block内thread的数量划分，这两者都是定义为 dim3 类型的变量。 dim3 可以视为包含三个无符号整数(x，y，z)的结构体变量，在定义时缺省值初始化为1，因此grid和block可以灵活地定义为一维、二维或三维结构。前面说到，每一个线程都会被分配唯一的线程ID。在上面的例子里，这个唯一的线程ID需要用两个内置变量blockIdx和threadIdx来唯一标识，它们都是 dim3 类型变量，其中blockIdx指明线程所属的block在grid中的索引，而threaIdx指明线程所在block中的索引，如图1中的Thread(3, 2)满足：

threadIdx.x = 3
threadIdx.y = 2
blockIdx.x = 1
blockIdx.y = 1

在这个例子里，Thread(1,1)的索引是第(1,1)个block内的第(3,2)个线程。在实际编程中，用户可以按需任意指定grid和block的数量，但需要注意这两者的数量是有限制的。在MUSA中，grid每个维度不能超过2147483647，grid中的每个block起码要有1个thread。同时，每个block的线程数量是有限制的，因为一个block的所有线程都需要驻留在同一个MP核心上，并且必须共享该MP核心的有限资源，比如register file和shared memory。在现有的GPU上，一个block最多可包含1024个线程。

通常来说，同一个问题可以有不同的grid/block大小配置(尽管在性能上可能会有一些差异)。一个具体的例子如下，以下代码将大小为MxN的两个矩阵A和B相加，并将结果存储到矩阵C中：

// Kernel definition
__global__ void MatAdd(float A[M][N], float B[M][N],
                       float C[M][N]) {
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[j][i] = A[j][i] + B[j][i];
}

int main() {
    ...
    // Kernel invocation with one block of M * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(M, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

此处执行的是block数量为1的kernel，其中唯一的一个block里面有MxN个线程。当MxN的大小超过单个block允许的最大值时，这样的线程划分就无法执行了，需要考虑其他做法。首先，在上面的代码中block数量由实参numBlocks指定为1，但同一个kernel是可以由很多个block并行执行的，这些block有同样的大小。因此在上面这个例子里，我们可以把矩阵分为若干个相同大小的子块，每个子块由一个与子块同样大小的block执行。此外，我们还需要计算当前线程在所有线程中的位置。这个位置可以用到前文中的线程ID，也就是blockIdx和threadIdx，以及另一个内置变量blockDim，也就是每个block的尺寸来求得：因为矩阵每个子块的大小都是相同的，分别求x和y方向上每个线程对应的矩阵索引(j, i)即可，代码如下：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N]) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[j][i] = A[j][i] + B[j][i];
}

int main() {
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

这里每个block的大小为16x16(共256个线程)，该大小可以任意指定(16x16是常见的划分)。上述例子中假设MxN矩阵的每个维度都可以由blcok中相同维度的线程数整除，但是在实际编程中更常见的情况是没法整除的，此时需要添加一些条件语句来判断是否属于可执行的thread，同时block数量要向上取整。所有block都是独立执行的，MUSA规定所有block可以以任何顺序执行(并行或串行)。这种独立性允许block以任意顺序在任意数量的内核上进行调度，从而使程序员能够编写随硬件核心数量扩展的代码。

同一block内的线程可以通过shared memory来共享数据，并通过同步block内所有线程的执行来协调内存的访问，以此进行协作。更准确地说，可以通过调用内联函数__syncthreads()来指定kernel中的同步点；__syncthreads()可以充当屏障，这意味着同一block内的所有线程都必须在该同步点等待，直到所有线程都运行到这一点才能继续往前执行。为了高效协作和执行，上述shared memory应该是每个MUSA核心附近的低延迟内存，并且__syncthreads()本身应该是轻量级的。

5.1.4. 内存层次结构

前面说到同一block内的线程可以利用shared memory进行数据共享。包括shared memory，在MUSA中存在多个层次的内存空间，如图2所示。在MUSA线程执行期间，每个线程都可以从多个内存空间访问数据。首先，每个线程都有私有的register和local memory；其次，每个block都有对该block中所有线程可见的共享内存，并且与该block具有相同的生命周期。最后，所有线程都可以访问相同的global memory。

此外，MUSA中还存在两个只读的内存空间可供所有线程访问：constant memory和texture memory。global memory、constant memory和texture memory三个内存空间针对不同的内存使用进行了优化，texture memory还可以为某些特定数据格式提供不同的寻址模式以及数据过滤，在实际使用中可以根据需求选择不同内存空间来实现任务需求。内存结构涉及到程序优化，这里暂不深入探讨。global memory、constant memory和texture memory内存空间在同一应用程序启动的kernel中保持不变。

5.1.5. 获取GPU硬件配置

对于任何硬件，想要用其实现高效的程序，都需要对其硬件信息有一定的了解。在进行MUSA编程前，良好的实践应该是首先检查GPU的硬件配置信息，再开始针对性的编程工作，这样才能做到有的放矢。但是实现这一目的有个常见问题，对于GPU来说，不同代的产品在硬件配置上可能有很大的差异，有时候用户可能无法从外观上获取相关信息或者需要针对多代GPU产品进行适配，这会是一件麻烦事。好在MUSA中为获取硬件信息提供了相关的查询API，用户当前使用的GPU硬件信息可以通过下面的代码来获得：

void getMUSADeviceProp() {
  int devID = 0;
  musaDeviceProp devProp;
  musaGetDeviceProperties(&devProp, devID);
  std::cout << "Device Name: " << devProp.name << std::endl;
  std::cout << "Device Warp Size: " << devProp.warpSize << std::endl;
  std::cout << "Device Register Per Block: " << devProp.regsPerBlock << std::endl;
  std::cout << "Device Register Per Multiprocessor: " << devProp.regsPerMultiprocessor << std::endl;
  std::cout << "Device Shared Memory Per Block: " << devProp.sharedMemPerBlock << std::endl;
  std::cout << "Device Shared Memory Per Multiprocessor: " << devProp.sharedMemPerMultiprocessor << std::endl;
  std::cout << "Device Max Threads Per Block: " << devProp.maxThreadsPerBlock << std::endl;
  std::cout << "Device Max Threads Per Multi Processor: " << devProp.maxThreadsPerMultiProcessor << std::endl;
  std::cout << "Device Total Global Memory: " << devProp.totalGlobalMem << std::endl;
  std::cout << "Device Total Const Memory: " << devProp.totalConstMem << std::endl;
  std::cout << "Device Max Threads Dim: " << devProp.maxThreadsDim[0] << ", " << devProp.maxThreadsDim[1] << ", " << devProp.maxThreadsDim[2] << std::endl;
  std::cout << "Device Max Grid Size: " << devProp.maxGridSize[0] << ", " << devProp.maxGridSize[1] << ", " << devProp.maxGridSize[2] << std::endl;
  std::cout << "Device L2 Cache Size: " << devProp.l2CacheSize << std::endl;
  return 0;
}
  // 输出如下:
  // Device Name: MTT S3000
  // Device Warp Size: 128
  // Device Register Per Block: 262144
  // Device Register Per Multiprocessor: 131072
  // Device Shared Memory Per Block: 28672
  // Device Shared Memory Per Multiprocessor: 28672
  // Device Max Threads Per Block: 1024
  // Device Max Threads Per Multi Processor: 6144
  // Device Total Global Memory: 34273656832
  // Device Total Const Memory: 163840
  // Device Max Threads Dim: 1024, 1024, 1024
  // Device Max Grid Size: 2147483647, 2147483647, 2147483647
  // Device L2 Cache Size: 524288

其中的devID用来指示是第几个设备，设备的ID从0开始递增。

5.1.6. MUSA编程实践

学习了上述MUSA编程基础相关的知识后，我们就可以开始一些简单的实践了。这里仍然用本节的第一个例子向量加法来举例，不同的是这里会考虑向量可能有任意长度。在开始之前，首先简单介绍下MUSA编程中内存管理相关的API。首先是在device上分配内存的musaMalloc函数。

__host__ __device__ musaError_t musaMalloc(void** devPtr, size_t size);

这个函数和C语言中的malloc十分类似，但是其操作对象是GPU的内存。musaMalloc可以在device上申请size字节的显存，其中devPtr是指向已分配内存的指针。从声明中可以看出，musaMalloc既可以在host端调用，也可以在device端调用。同样的，类似c中的free，要释放已分配的内存，应该使用musaFree函数。需要说明的是musaMalloc申请出来的显存，其内容不保证预先清零，与之相近行为的接口是muMemCreate。

__host__ __device__ musaError_t musaFree(void* devPtr);

另外一个重要的API是负责host和device之间数据传输的函数musaMemcpy：

__host__ musaError_t musaMemcpy(void* dst, const void* src, size_t count, musaMemcpyKind kind);

其中src是指向源数据的指针；dst则是指向目标数据的指针；count是需要传输的数据大小，单位是字节；最后一个参数kind是枚举变量，用于控制数据传输的方向，包括musaMemcpyHostToHost、musaMemcpyHostToDevice、musaMemcpyDeviceToHost和musaMemcpyDeviceToDevice，比如musaMemcpyHostToDevice意味着将host上的数据复制到device上。musaMemcpy只能在host端调用。

现在我们来实现该向量加法，目标是计算向量z = x + y。这里的grid和block都设为一维，首先定义kernel如下：

__global__ void vecAdd(const float* x, const float * y, float* z, int n) {
    // get index
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // stride
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
    {
        z[i] = x[i] + y[i];
    }
}

因为grid和block都是一维的，所以stride等于所有block的长度之和，即整个grid的线程数t。当向量长度n <= t时，代码中的循环只执行一次(或者不执行)。现在考虑向量长度n > t的情况，此时可以在每个线程中实现多个元素(元素总数/线程总数)的加法。这被称为grid-stride loop方法。当前n的数量很大时，我们可以将grid的线程数t设为刚好与硬件的MUSA核心数量相等，这样可以充分利用所有硬件核心的性能，代价是某一部分核心在i < n时可能会“空跑”；这个例子里也可以将grid线程数量设为其他值，除了性能差别外不影响最后结果的一致性。以下是具体的代码实现：

#include <musa.h>

int main() {
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // host memory
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);

    // initiate host data x and y
    for (int i = 0; i < N; ++i) {
        x[i] = i + 1.0f;
        y[i] = i + 2.0f;
    }

    // device memory
    float *d_x, *d_y, *d_z;
    musaMalloc((void**)&d_x, nBytes);
    musaMalloc((void**)&d_y, nBytes);
    musaMalloc((void**)&d_z, nBytes);

    // copy host data to device
    musaMemcpy((void*)d_x, (void*)x, nBytes, musaMemcpyHostToDevice);
    musaMemcpy((void*)d_y, (void*)y, nBytes, musaMemcpyHostToDevice);

    // define grid and block layout
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);

    // execute kernel
    vecAdd<<<gridSize, blockSize>>>(d_x, d_y, d_z, N);

    // copy device data back to host
    musaMemcpy((void*)z, (void*)d_z, nBytes, musaMemcpyDeviceToHost);

    // free device memory
    musaFree(d_x);
    musaFree(d_y);
    musaFree(d_z);

    // free host memory
    free(x);
    free(y);
    free(z);
    return 0;
}

代码中向量的长度n设为1<<20，如果取block大小为256，那么grid的大小就是n/256=4096。此时kernel的线程层级结构如下图所示，对于其中的黄色方格所在的线程，其index是515。

需要注意的是kernel执行是与host异步的，这里使用musaMemcpy()函数会隐式地保证device和host同步之后再执行数据拷贝；这里也可以用显式的同步函数cudaDeviceSynchronize()。这样操作之后，下一步的代码才可以正确访问kernel计算的结果。

5.2. MUSA C++ API

5.2.1. 函数执行空间说明符

函数执行空间说明符表示函数是在主机上执行还是在设备上执行，以及它是可从主机调用还是可从设备调用。

5.2.1.1. global_

__global__ 执行空间说明符将函数声明为内核。这样的函数：

• 在设备上执行
• 可以从主机调用

__global__ 函数必须具有void返回类型，并且不能是类的成员。

对 __global__ 函数的任何调用都必须指定其执行配置。

对 __global__ 函数的调用是异步的，这意味着它在设备完成执行之前返回。

5.2.1.2. device

__device__ 执行空间说明符声明的函数：

• 在设备上执行
• 仅能从设备上调用

__global__ 和 __device__ 执行空间说明符不能一起使用。

5.2.1.3. host

__host__ 执行空间说明符声明的函数：

• 在主机上执行
• 仅能从主机上调用

只使用 __host__ 执行空间说明符来声明函数，等价于不使用 __host__、__device__ 或 __global__ 的任何一个执行空间描述符来声明函数；在这两种情况下，函数都只是为主机编译的。

__global__ 和 __host__ 执行空间说明符不能一起使用。

但是，__device__ 和 __host__ 执行空间说明符可以一起使用，在这种情况下，为主机和设备编译函数。应用程序兼容性中引入的 __MUSA_ARCH__ 宏可用于区分主机和设备之间的代码路径：

__host__ __device__ void func()
{
    #if defined(__MUSA_ARCH__)
        // device code path
    #else
        // host code path
    #endif
}

5.2.1.4. 未定义的行为

“交叉执行空间”调用在以下情况下具有未定义的行为：

__MUSA_ARCH__ 有定义时从 __global__、__device__ 或 __host__ __device__ 函数内部对 __host__ 函数的调用。

__MUSA_ARCH__ 未定义时从 __host__ 函数内部对 __device__ 函数的调用。

5.2.1.5. noinline 和 forceinline

编译器在适当的时候内联任何 __device__ 函数。

如果可能的话，函数限定符 __noinline__ 可以用作编译器不内联函数的提示。

函数限定符 __forceinline__ 可用于强制编译器内联函数。

函数限定符 __noinline__ 和 __forceinline__ 不能一起使用，并且这两个函数限定符都不能应用于内联函数。

5.2.2. 变量存储空间说明符

变量存储空间说明符表示变量在设备上的内存位置。

在设备代码中声明的一个自动变量，没有本节中描述的任何 __device__ 、__shared__ 和 __constant__ 存储空间说明符，通常驻留在寄存器中。然而，在某些情况下，编译器可能会选择将其放置在本地存储器中。

5.2.2.1. __device__

__device__ 存储空间说明符声明一个驻留在设备上的变量。

在接下来的三个部分中定义的其他存储空间说明符中，最多可以有一个与 __device__ 一起使用，以进一步表示变量属于哪个存储空间。如果它们都不存在，则变量：

• 位于全局存储空间中,
• 具有创建 MUSA 上下文的生存期，
• 每个设备具有不同的对象，
• 可从网格内的所有线程以及主机端通过运行时库（musaGetSymbolAddress() / musaGetSymbolSize() / musaMemcpyToSymbol() / musaMemcpyFromSymbol()）进行访问。

5.2.2.2. __constant__

__constant__ 存储空间说明符（可与 __device__ 一起使用）声明一个变量，该变量：

• 驻留在常量存储空间中，
• 具有创建 MUSA 上下文的生存期，
• 每个设备具有不同的对象，
• 可从网格内的所有线程以及主机通过运行时库（musaGetSymbolAddress() / musaGetSymbolSize() / musaMemcpyToSymbol() / musaMemcpyFromSymbol()）进行访问。

5.2.2.3. __shared__

__shared__ 存储空间说明符（可与 __device__ 一起使用）声明一个变量，该变量：

• 驻留在线程块的共享内存空间中，
• 具有线程块的生命周期，
• 每个块具有不同的对象，
• 仅可从块内的所有线程访问，
• 没有固定地址。

将共享内存中的变量声明为外部数组时，例如

extern __shared__ float shared[];

数组的大小在启动时确定（请参阅执行配置）。以这种方式声明的所有变量都从共享内存中的同一地址开始，因此数组中变量的布局必须通过偏移量进行显式管理。例如，如果想要在动态分配的共享内存中声明：

short array0[128];
float array1[64];
int   array2[256];

可以通过以下方式声明和初始化数组：

extern __shared__ float array[];
__device__ void func()      // __device__ or __global__ function
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[128];
    int*   array2 =   (int*)&array1[64];
}

请注意，指针需要与它们所指向的类型对齐，因此以下代码不起作用，例如，因为 array1 没有 4 个字节对齐。

extern __shared__ float array[];
__device__ void func()      // __device__ or __global__ function
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[127];
}

5.2.2.4. __restrict__

mcc 通过 __restrict__ 关键字支持受限指针。

C99 中引入了受限指针，以缓解 C 类型语言中存在的别名问题，并抑制从代码重新排序到公共子表达式消除的所有优化。

以下是一个涉及别名问题的示例，其中使用受限指针可以帮助编译器减少指令数量：

void foo(const float* a,
         const float* b,
         float* c)
{
    c[0] = a[0] * b[0];
    c[1] = a[0] * b[0];
    c[2] = a[0] * b[0] * a[1];
    c[3] = a[0] * a[1];
    c[4] = a[0] * b[0];
    c[5] = b[0];
    ...
}

在 C 类型语言中，指针 a、b 和 c 可能是别名，因此任何直写 c 都可以修改 a 或 b 的元素。这意味着，为了保证函数的正确性，编译器不能将 a[0] 和 b[0] 加载到寄存器中，将它们相乘，并将结果存储到 c[0] 和 c[1]，因为 a[0] 与 c[0] 有可能实际上是同一个位置。因此编译器无法利用公共子表达式。同样，编译器不能将 c[4] 的计算重新排序为与 c[0] 和 c[1] 的计算附近，因为之前对 c[3] 的写入可能会改变对 c[4]计算的输入。

通过将 a、b 和 c 设置为受限指针，程序员向编译器断言指针实际上没有别名，这意味着通过 c 写入永远不会覆盖 a 或 b 的元素。这将改变函数原型，如下所示：

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c);

请注意，需要对所有指针参数进行限制，编译器优化器才有可能获得好处。添加了 __restrict__ 关键字后，编译器现在可以随意重新排序和消除公共子表达式，同时保留与抽象执行模型相同的功能：

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c)
{
    float t0 = a[0];
    float t1 = b[0];
    float t2 = t0 * t1;
    float t3 = a[1];
    c[0] = t2;
    c[1] = t2;
    c[4] = t2;
    c[2] = t2 * t3;
    c[3] = t0 * t3;
    c[5] = t1;
    ...
}

这里的效果是减少了内存访问次数和减少了计算次数。由于“缓存”的加载和常见的子表达式，寄存器压力的增加平衡了这一点。

由于寄存器压力是许多 MUSA 代码中的一个关键问题，由于占用率降低，使用受限指针可能会对 MUSA 代码的性能产生负面影响。

5.2.3. 内置变量类型

5.2.3.1. char, short, int, long, longlong, float, double

这些是从基本整数和浮点类型派生而来的向量类型。它们是结构体，并且可以分别通过字段 x、y、z 和 w 访问第一、第二、第三和第四分量。它们都带有一个形式为 `make_

` 的构造函数；例如

int2 make_int2(int x, int y);

创建类型为 int2 的向量，其值为 (x, y)。

下表详细介绍了矢量类型的对齐要求。

类型	对齐要求
char1, uchar1	1
char2, uchar2	2
char3, uchar3	1
char4, uchar4	4
short1, ushort1	2
short2, ushort2	4
short3, ushort3	2
short4, ushort4	8
int1, uint1	4
int2, uint2	8
int3, uint3	4
int4, uint4	16
long1, ulong1	如果 sizeof(long) 等于 sizeof(int) 为 4 否则为 8
long2, ulong2	如果 sizeof(long) 等于 sizeof(int) 为 8 否则为 16
long3, ulong3	如果 sizeof(long) 等于 sizeof(int) 为 4 否则为 8
long4, ulong4	16
longlong1, ulonglong1	8
longlong2, ulonglong2	16
longlong3, ulonglong3	8
longlong4, ulonglong4	16
float1	4
float2	8
float3	4
float4	16
double1	8
double2	16
double3	8
double4	16

5.2.3.2. dim3

此类型是基于 uint3 的整数向量类型，用于指定维度。当定义一个 dim3 类型的变量时，任何未指定的组件被初始化为 1。

5.2.4. 内置变量

内置变量指定网格和块尺寸以及块和线程索引。它们仅在设备上执行的函数中有效。

5.2.4.1. gridDim

此变量是 dim3 类型（参见 dim3）的，用于指定网格的尺寸。例如，如果网格的尺寸为 (x, y, z)，则可以通过 gridDim.x、gridDim.y 和 gridDim.z 访问 x、y 和 z。

5.2.4.2. blockIdx

此变量是 uint3 类型（参见 char, short, int, long, longlong, float, double）的，用于指定网格中块的索引。

5.2.4.3. blockDim

此变量是 dim3 类型（参见 dim3）的，用于指定块的尺寸。

5.2.4.4. threadIdx

此变量是 uint3 类型（参见 char, short, int, long, longlong, float, double）的，用于指定块中线程的索引。

5.2.4.5. warpSize

此变量是 int 类型的，用于指定线程束的大小。

5.2.5. 同步函数

void __syncthreads();

在线程块中的所有线程都达到这一点，并且这些线程在 __syncthreads() 之前进行的所有全局和共享内存访问对块中的所有线程都可见之前等待。

__syncthreads() 用于协调同一块的线程之间的通信。当块中的某些线程访问共享内存或全局内存中的相同地址时，其中一些内存访问可能存在写后读、读后写或写后写风险。可以通过在这些访问之间同步线程来避免这些访存风险。

__syncthreads() 在条件代码中是允许的，但前提是条件在整个线程块中的计算方式相同，否则代码执行可能会挂起或产生意外的副作用。

__syncthreads() 支持下面描述的三种变体。

void __syncthreads_count(int predicate);

基本功能与 __syncthreads() 功能相同，附加功能是它会针对块的所有线程计算 predicate，并返回计算结果为非零的线程数。

int __syncthreads_and(int predicate);

基本功能与 __syncthreads() 功能相同，附加功能是它会针对块的所有线程计算 predicate，并当且仅当所有线程的计算结果都为非零时返回非零值。

int __syncthreads_or(int predicate);

基本功能与 __syncthreads() 功能相同，附加功能是它会针对块的所有线程计算 predicate，并当且仅当至少一个线程的计算结果为非零时返回非零值。

void __syncwarp(unsigned mask=0xffffffff);

将导致执行线程等待，直到掩码中覆盖的同一线程束内所有线程都执行了一次 __syncwarp()（使用相同的掩码），然后再恢复执行。每个调用线程必须在掩码中设置自己的位，并且掩码中覆盖的所有未退出线程必须使用相同的掩码执行相应的操作，否则结果是未定义的。

执行 __syncwarp() 保证参与栅栏同步的线程之间的内存顺序。因此，线程束中希望通过内存进行通信的线程可以存储值到内存中，执行 __syncwarp() ，然后安全地读取同一线程束中其他线程存储的值。

5.2.6. 原子函数

原子函数对驻留在全局内存或共享内存中的一个 32 位或 64 位数据执行读-修改-写原子操作。例如，atomicAdd() 读取全局内存或共享内存中某个地址的一个字大小的值，向其增加一个数值，然后将结果写回同一地址。原子函数只能在设备函数中使用。

请注意，任何原子操作都可以基于 atomicCAS()（比较和交换）实现。例如，双精度浮点数操作可以按如下方式实现：

  __device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}

5.2.6.1 算术原子函数

5.2.6.1.1. atomicAdd()

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);

读取位于全局或共享内存中地址 address 存储的 16 位、32 位或 64 位 old 值，计算 (old + val)，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.2. atomicSub()

int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address,
                       unsigned int val);

读取位于全局或共享内存中地址 address 存储的 32 位 old 值，计算 (old - val)，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.3. atomicExch()

int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address,
                        unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address,
                                  unsigned long long int val);
float atomicExch(float* address, float val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，将 val 存储到同一地址的内存中。这两个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.4. atomicMin()

int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMin(long long int* address,
                                long long int val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，计算 old 和 val 中的最小值，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.5. atomicMax()

int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMax(long long int* address,
                                 long long int val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，计算 old 和 val 中的最大值，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.6. atomicInc()

unsigned int atomicInc(unsigned int* address,
                       unsigned int val);

读取位于全局或共享内存中地址 address 存储的 32 位 old 值，计算 ((old >= val) ? 0 : (old+1))，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.7. atomicDec()

unsigned int atomicDec(unsigned int* address,
                       unsigned int val);

读取位于全局或共享内存中地址 address 存储的 32 位 old 值，计算 (((old == 0) || (old > val)) ? val : (old-1))，并将结果存储回同一地址的内存。这三个操作在一个原子事务中执行。该函数返回 old。

5.2.6.1.8. atomicCAS()

int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address,
                       unsigned int compare,
                       unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,
                                 unsigned long long int compare,
                                 unsigned long long int val);
unsigned short int atomicCAS(unsigned short int *address,
                             unsigned short int compare,
                             unsigned short int val);

读取位于全局或共享内存中地址 address 存储的 16 位、32 位或 64 位 old 值，计算 (old == compare ? val : old)，并将结果存储到同一地址的内存中。这两个操作在一个原子事务中执行。该函数返回 old（比较与交换）。

5.2.6.2. 位原子操作

5.2.6.2.1. atomicAnd()

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，计算 (old & val)，并将结果存储回同一地址的内存。这两个操作在一个原子事务中执行。该函数返回 old。

5.2.6.2.2. atomicOr()

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
                      unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
                                unsigned long long int val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，计算 (old | val)，并将结果存储回同一地址的内存。这两个操作在一个原子事务中执行。该函数返回 old。

5.2.6.2.3. atomicXor()

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址 address 存储的 32 位或 64 位 old 值，计算 (old ^ val)，并将结果存储回同一地址的内存。这两个操作在一个原子事务中执行。该函数返回 old。

5.2.7. 线程束投票函数

int __all_sync(unsigned mask, int predicate);
int __any_sync(unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);
unsigned __activemask();

线程束投票函数允许给定线程束的线程执行归约和广播操作。这些函数将来自线程束中中每个线程的整数类型 predicate 作为输入，并将这些值与零进行比较。比较的结果通过以下方式之一在线程束中的活动线程之间组合（归约），将单个返回值广播到每个参与线程：

• __all_sync(unsigned mask, predicate)：评估 mask 所包含的所有未退出线程的 predicate 值，当且仅当所有线程的计算结果为非零时返回非零。
• __any_sync(unsigned mask, predicate)：评估 mask 所包含的所有未退出线程的 predicate 值，当且仅当至少一个线程的计算结果为非零时返回非零。
• __ballot_sync(unsigned mask, predicate)：评估 mask 所包含的所有未退出线程的 predicate 值，并将结果作为一个 32 位整数返回。每个线程的返回值都对应整数一位，如果线程未退出且 predicate 值非零则该位为 1，否则为 0。
• __activemask()：返回一个 32 位整数，每个线程都对应整数的一位。如果线程未退出则该位为 1，否则为 0。

对于 __all_sync、__any_sync 和 __ballot_sync，必须传递一个掩码，用于指定参与调用的线程。必须为每个参与线程设置一个表示线程束通道索引的位，以确保它们在硬件执行内部函数之前正确收敛。每个调用线程必须在掩码中设置自己的位，并且掩码中命名的所有非退出线程必须使用相同的掩码执行相同的内部函数，否则结果是未定义的。

这些内置函数并不意味着栅栏同步。它们不保证任何访存顺序。

5.2.8. 线程束值交换函数

__shfl_sync、__shfl_up_sync、__shfl_down_sync 以及 __shfl_xor_sync 在线程束内的线程间交换一个值。

5.2.8.1. 简介

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

T 可以是 int、unsigned int、long、unsigned long、long long、 unsigned long long、float 或 double。

5.2.8.2. 详细描述

__shfl_sync() 内置函数允许在线程束中的线程之间交换变量，而无需使用共享内存。对于线程束内的所有活动（且被 mask 包含的）线程，交换是同时发生的，根据类型，每个线程移动 4 或 8 个字节的数据。

线程束内的线程称为通道，其索引可能介于 0 和 warpSize-1（含）之间。支持以下四种源通道寻址模式：

• __shfl_sync() 从被索引的通道中直接拷贝
• __shfl_up_sync() 从相较于调用线程索引更低的通道中拷贝
• __shfl_down_sync() 从相较于调用线程索引更高的通道中拷贝
• __shfl_xor_sync() 从调用线程索引的异或结果通道中拷贝

线程只能从主动参与 __shfl_sync() 命令的另一个线程读取数据。如果目标线程处于非活动状态，则返回的值未定义。

所有 __shfl_sync() 内置函数都采用一个可选 width 参数，该参数会改变内置函数的行为。width 的值必须是 2 的幂；如果宽度不是 2 的幂，或者是一个大于 warpSize 的数值，则结果未定义。

__shfl_sync() 返回由 srcLane 给定的索引对应的线程所持有的 var 值。如果宽度小于 warpSize，则线程束内的每个子部分都表现为起始逻辑通道索引为 0 的单独实体。如果 srcLane 超出范围 [0:width-1]，则返回的值对应于 srcLane modulo width（即在同一子部分内）对应的线程持有的 var 值。

__shfl_up_sync() 通过从调用者的通道索引中减去 delta 来计算源通道索引。返回由生成的通道索引对应的线程所持有的 var 值：实际上，var 在线程束内向上移动了 delta 个通道。如果 width 小于 warpSize，则线程束的每个子部分都表现为起始逻辑通道索引为 0 的单独实体。源通道索引不会环绕 warpSize 的值，因此实际上 delta 较低的通道将保持不变。

__shfl_down_sync()通过增加 delta 到调用者的通道索引来计算源通道索引。返回生成的通道索引对应的线程所持有的 var 值：实际上，var 在线程束内向下移动了 delta 个通道。如果 width 小于 warpSize，则线程束的每个子部分都表现为起始逻辑通道索引为 0 的单独实体。源通道索引不会环绕 warpSize 值，因此实际上 delta 较高的通道将保持不变。

__shfl_xor_sync() 通过对调用者的通道索引与 laneMask 执行按位异或来计算源通道索引，并返回由生成的通道索引对应的线程所持有的 var 值。如果 width 小于 warpSize，则每组 wdith 个连续线程都能够访问较早线程组中的元素，但是，如果它们尝试访问后面线程组中的元素，则将返回它们自己的 var 值。此模式实现蝴蝶寻址模式，例如用于树归约和广播。

线程束值交换内部函数采用掩码指示参与调用的线程。必须为每个参与线程设置一个表示线程通道索引的位，以确保在硬件执行这些内置函数之前正确收敛到它们。每个调用线程必须在掩码中设置自己的位，并且掩码中包含的所有非退出线程必须使用相同的掩码执行相同的内部函数，否则结果是未定义的。

这些内置函数并不意味着栅栏同步。它们不保证任何访存顺序。

5.2.8.3. 示例

5.2.8.3.1. 将单个值广播到线程束其他线程

#include <stdio.h>

__global__ void bcast(int arg) {
    int laneId = threadIdx.x & 0x1f;
    int value;
    if (laneId == 0)        // Note unused variable for
        value = arg;        // all threads except lane 0
    value = __shfl_sync(0xffffffff, value, 0);   // Synchronize all threads in warp, and get "value" from lane 0
    if (value != arg)
        printf("Thread %d failed.\n", threadIdx.x);
}

int main() {
    bcast<<< 1, 32 >>>(1234);
    cudaDeviceSynchronize();

    return 0;
}

5.2.8.3.2. 在以 8 个线程为单位的组中计算包含前缀和

#include <stdio.h>

__global__ void scan4() {
    int laneId = threadIdx.x & 0x1f;
    // Seed sample starting value (inverse of lane ID)
    int value = 31 - laneId;

    // Loop to accumulate scan within my partition.
    // Scan requires log2(n) == 3 steps for 8 threads
    // It works by an accumulated sum up the warp
    // by 1, 2, 4, 8 etc. steps.
    for (int i=1; i<=4; i*=2) {
        // We do the __shfl_sync unconditionally so that we
        // can read even from threads which won't do a
        // sum, and then conditionally assign the result.
        int n = __shfl_up_sync(0xffffffff, value, i, 8);
        if ((laneId & 7) >= i)
            value += n;
    }

    printf("Thread %d final value = %d\n", threadIdx.x, value);
}

int main() {
    scan4<<< 1, 32 >>>();
    cudaDeviceSynchronize();

    return 0;
}

5.2.8.3.3. 在线程束内计算归约

#include <stdio.h>

__global__ void warpReduce() {
    int laneId = threadIdx.x & 0x1f;
    // Seed starting value as inverse lane ID
    int value = 31 - laneId;

    // Use XOR mode to perform butterfly reduction
    for (int i=16; i>=1; i/=2)
        value += __shfl_xor_sync(0xffffffff, value, i, 32);

    // "value" now contains the sum across all threads
    printf("Thread %d final value = %d\n", threadIdx.x, value);
}

int main() {
    warpReduce<<< 1, 32 >>>();
    cudaDeviceSynchronize();

    return 0;
}

5.2.9. 格式化输出

int printf(const char *format[, arg, ...]);

将格式化的输出从 kernel 打印到主机端输出流。

kernel 内 printf() 函数的行为方式与标准 C 库 printf() 函数类似，用户可参阅主机系统的手册页以获取完整的行为描述。实质上，传入的字符串将输出到主机上的流，并在遇到格式说明符的任何地方从参数列表进行替换。下面列出了支持的格式说明符。

该命令像任何其他设备端函数一样执行：每线程和在调用线程的上下文中。这意味着在拥有多线程的 kernel 中，每个线程都将使用该线程的数据执行直接调用。然后，输出字符串的多个版本将出现在主机流中，每个对应于一个线程执行了一次 printf()。

如果只需要一个输出字符串，则由程序员将输出限制为单个线程（有关说明性示例，请参阅示例）。

与返回打印字符数的 C 标准 printf() 不同，MUSA 返回解析的参数数。如果格式字符串后面没有参数，则返回 0。如果格式字符串为 NULL，则返回 -1。如果发生内部错误，则返回 -2。

5.2.9.1. 格式说明符

与标准 printf() 一样，格式说明符采用以下形式：%[flags][width][.precision][size]type

支持以下字段（有关所有行为的完整说明，请参阅相关文档）：

• 标志：'#' ' ' '0' '+' '-'
• 宽度：'*' '0-9'
• 精度：'0-9'
• 大小：'h' 'l' 'll'
• 类型："%cdiouxXpeEfgGaAs"

请注意，MUSA 将接受标志、宽度、精度、大小和类型的任何组合，无论它们是否总体上形成有效的格式说明符。换句话说，“%hd” 将被接受，printf() 将期望在参数列表中的相应位置有一个双精度变量。

5.2.9.2. 示例

以下代码：

#include <stdio.h>

__global__ void helloMUSA(float f)
{
    printf("Hello thread %d, f=%f\n", threadIdx.x, f);
}

int main()
{
    helloMUSA<<<1, 5>>>(1.2345f);
    musaDeviceSynchronize();
    return 0;
}

将以下输出打印到标准输出：

Hello thread 2, f=1.2345
Hello thread 1, f=1.2345
Hello thread 4, f=1.2345
Hello thread 0, f=1.2345
Hello thread 3, f=1.2345

请注意每个线程都遇到了 printf() 命令，因此输出行数与网格中启动的线程数一样多。正如预期的那样，全局值（即 float f）在所有线程之间是通用的，而本地值（即 threadIdx.x）在每个线程之间是不同的。

以下代码：

#include <stdio.h>

__global__ void helloMUSA(float f)
{
    if (threadIdx.x == 0)
        printf("Hello thread %d, f=%f\n", threadIdx.x, f) ;
}

int main()
{
    helloMUSA<<<1, 5>>>(1.2345f);
    musaDeviceSynchronize();
    return 0;
}

将以下输出打印到标准输出：

Hello thread 0, f=1.2345

不言而喻，if() 语句限制了哪些线程将调用 printf，因此只看到一行输出。

5.2.10. 执行配置

对 __global__ 函数的任何调用都必须指定该调用的执行配置。执行配置定义了将用于在设备上执行函数的网格和块的维度，以及关联的流（有关流的说明，请参阅 MUSA 运行时）。

通过在函数名称和带括号的参数列表之间插入形如<<< Dg, Db, Ns, S >>> 的表达式来指定执行配置，其中：

• Dg 为 dim3 类型（参见 dim3），指定网格的尺寸和大小，这样 Dg.x * Dg.y * Dg.z 就等于要启动的线程块数量；
• Db 为 dim3 类型（请参阅 dim3），指定每个线程块的尺寸和大小，这样 Db.x * Db.y * Db.z 就等于每个块的线程数；
• Ns 为 size_t 类型，指定除了静态分配的共享内存之外，每个线程块为此次调用动态分配的共享内存中字节数；如 __shared__ 中所述，此动态分配的内存由声明为外部数组的任何变量使用；Ns 是一个可选参数，默认为 0;
• S 为 cudaStream_t 类型并指定关联的流；S 是一个默认为 0 的可选参数。

例如，声明为

__global__ void Func(float* parameter);

必须用如下形式调用：

Func<<< Dg, Db, Ns >>>(parameter);

执行配置的参数在实际函数参数之前计算。

如果 Dg 或 Db 大于计算功能中指定的设备允许的最大大小，或者 Ns 大于设备上可用的最大共享内存量减去静态分配所需的共享内存量，则函数调用将失败。

5.2.11. 循环展开

默认情况下，编译器展开具有已知行程计数的小循环。但是，#pragma unroll 指令可用于控制任何给定循环的展开。它必须放在紧靠循环之前，并且仅适用于该循环。它后面可选地跟一个整数常量表达式。如果常量表达式不存在，且循环的行程计数恒定，则循环将完全展开。如果常量表达式的计算结果为 1，编译器将不会展开循环。如果常量表达式的计算结果为非正整数或大于 int 类型可表示的最大值，则将忽略该语句。

示例：

struct S1_t { static const int value = 4; };
template <int X, typename T2>
__device__ void foo(int *p1, int *p2) {

// no argument specified, loop will be completely unrolled
#pragma unroll
for (int i = 0; i < 12; ++i)
  p1[i] += p2[i]*2;

// unroll value = 8
#pragma unroll (X+1)
for (int i = 0; i < 12; ++i)
  p1[i] += p2[i]*4;

// unroll value = 1, loop unrolling disabled
#pragma unroll 1
for (int i = 0; i < 12; ++i)
  p1[i] += p2[i]*8;

// unroll value = 4
#pragma unroll (T2::value)
for (int i = 0; i < 12; ++i)
  p1[i] += p2[i]*16;
}

__global__ void bar(int *p1, int *p2) {
foo<7, S1_t>(p1, p2);
}

5.3. mcc编译器基础使用

mcc随MUSA Toolkit安装到用户环境后即可使用。mcc将尝试在MUSA Toolkit默认安装路径“/usr/local/musa”下检测是否正确安装MUSA Toolkit。若用户将MUSA Toolkit安装至非默认路径下，mcc支持通过命令行选项“--musa-path”指定MUSA Toolkit的安装位置。

当前mcc仅支持Linux平台，Windows版本的编译器将在后续版本中发布。

5.3.1. 基础编译命令

mcc基于clang-14与llvm开发，基础的使用方式：包括链接库、库文件路径、头文件路径等选项与开源clang保持一致。我们以MUSA源码文件：axpy.mu为例，使用如下命令将它编译成可执行文件：

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

其中：

• axpy.mu: 输入的源文件；
• -lmusart：链接MUSA runtime动态链接库，与clang++编译普通c++源码文件时，指定依赖的动态链接库的命令行选项格式保持一致；
• -L/usr/local/musa/lib：指定MUSA runtime所在的路径，与clang++编译普通c++源码时，指定动态链接库路径的命令行选项格式保持一致；
• -o axpy：表示指定输出的可执行文件名为“axpy"

5.3.2. mcc编译器编译选项

mcc继承了大部分clang++的命令行选项，以下为常用或mcc扩展的命令行选项说明：

5.3.2.1. 帮助文档类选项

--help/-help

显示命令行选项帮助文档。

--help-hidden

显示默认隐藏的命令行选项的帮助文档。

-v

显示编译过程中的详细编译命令。

--version

打印版本信息。

5.3.2.2. MUSA Toolkit安装路径相关选项

`--musa-path=

`

指定MUSA Toolkit的安装路径。

`--musa-inc-path=

`

（如需）单独指定MUSA Toolkit中头文件的路径；默认路径为：`

/include`

`--musa-libdevice-path=

`

（如需）单独指定MUSA Toolkit中libdevice.bc所在的路径；默认路径为：`

/mtgpu/bitcode/`

5.3.2.3. 其他选项

`-o

`

指定编译产物的文件名。

-x musa

mcc默认通过识别源文件后缀名为“.mu”时，将源码识别为MUSA源码，而当输入文件为其他后缀时，可以通过该选项指定非“.mu”后缀的源代码为MUSA源码。该组选项必须加在源文件名在命令行中的相对位置之前。

完整的mcc编译器命令行选项请参考编译器用户手册。