跳到主要内容

MCC 用户手册

本文档描述了如何使用 mcc 编译 MUSA 代码，并提供了一些关于 mcc 实现的详细信息。

本文档假设您对 CUDA 编程语言有基本的了解。有关 CUDA 编程的信息可以在 CUDA 编程指南中找到。

概览​

MUSA是一种类似于 CUDA 的编程语言，用于在 MTGPU 设备上进行并发编程。

先决条件​

mcc 目前支持构建与 CUDA 11.4 兼容的 MUSA。在构建 MUSA 代码之前，您需要安装 MUSA Toolkit。mcc 将尝试在默认路径 “/usr/local/musa” 下找到 MUSA 安装。用户也可以通过 “--musa-path” 选项指定 MUSA Toolkit 位置的不同路径。

MUSA 编译目前仅在 Linux 上支持。Windows 上的编译器即将推出。

MCC的目的​

MUSA应用程序由传统的 C++ 主机代码和 GPU 设备函数混合而成，这使得编译过程涉及每个 MUSA 源文件的多个编译和链接步骤。 mcc，MUSA 编程语言编译器驱动程序，有助于隐藏构建源代码的这些阶段的复杂性。

编译阶段​

大多数情况下，mcc 中的 MUSA 编译与传统 C++ 类似，使用 clang++：

  1. 预处理源代码，通过宏扩展和头文件包含来处理宏和包含指令。
  2. 编译为 llvm ir，通过一系列有组织的转换和优化传递。
  3. 代码生成：降低到目标特定代码并生成对象文件。
  4. 链接库以生成可执行二进制文件，包括静态/共享库和独立可执行文件。

不同之处在于 MUSA 代码与主机端部分和设备端部分混合。基本上，mcc 将编译源代码两次以生成设备代码对象和主机代码对象，然后尝试将它们组合在一起。因此，上述阶段至少会为设备端代码和主机端部分重复两次。

详细来说，在代码生成之前，设备代码对象的二进制值将被分配给全局常量变量，当编译主机端代码时。主机端代码的阶段与传统 C++ 几乎相同。

当内核被发射到 GPU 设备时，MUSA 运行时和驱动程序将帮助提取设备代码对象，这超出了本文档的范围。

编译轨迹​

由mcc驱动的编译遵循以下轨迹： 编译轨迹

支持的输入文件后缀​

由mcc驱动的编译的每个阶段都接受特定的输入文件。mcc将通过文件后缀和编译选项来识别它们。

输入文件后缀描述
.mu一个 MUSA 源文件，包含主机代码和设备函数。
.c一个 C 源文件。
.cc, .cxx, .cpp一个 C++ 源文件。
.ll一个 LLVM IR 文件。
.o, .obj一个对象文件。
.a一个库文件。
.so一个共享库文件。

其他文件后缀由 clang/clang++ 支持。

编译器命令行选项​

mcc是MUSA语言的编译器驱动程序，受 llvm-project 14.0.0 启发。大多数编译器选项与 Clang 和 LLVM 工具链一样有效。此外，它还为 MUSA 语言和 MTGPU 目标提供了特定选项。

帮助信息选项​

--help/-help​

显示可用选项的帮助。

--help-hidden​

显示隐藏选项的帮助。

-v​

显示子命令的详细输出。

--version​

打印版本信息。

输入文件和规范选项​

-o [file]​

定义输出文件的名称。

-x [language]​

指定输入文件的语言（例如 -x musa），后续输入文件将被识别为具有类型。

-D [macro]=[value]​

定义[macro]为[value]（如果省略[value]则为1）。

路径规范选项​

--musa-path= value​

指定MUSA Toolkit安装路径。

--musa-path-ignore-env​

忽略环境变量以检测 MUSA Toolkit 安装。

--musa-inc-path= value​

指定 MUSA Toolkit 中头文件的路径，默认路径是 musa-path/include。

--musa-libdevice-path= value​

指定 MUSA Toolkit中libdevice.bc 的路径，默认路径是 musa-path/mtgpu/bitcode/。

--musa-rt-lib-path= value​

指定MUSA运行时库路径。

包含和链接选项​

-I dir​

将目录添加到包含搜索路径列表的末尾。

-include file​

在解析之前包含文件。

-l lib-name​

与库链接。

-L dir​

将目录添加到库搜索路径。

-nogpuinc​

不添加 CUDA/HIP/MUSA 的包含路径，并且不包括默认的 CUDA/HIP/MUSA 包装器头文件。

-nogpulib​

不链接 CUDA/HIP/MUSA 设备编译的设备库。

编译阶段控制选项​

-c​

仅运行预处理、编译和汇编步骤。

-emit-llvm​

使用LLVM表示汇编器和对象文件。

--musa-compile-host-device​

为主机和设备编译 MUSA 代码（默认）。对非 MUSA 编译无影响。

--musa-device-only​

仅为设备编译 MUSA 代码。

--musa-host-only​

仅为主机编译 MUSA 代码。对非 MUSA 编译无影响。

--offload-arch= value​

MUSA 卸载目标 ID 以设备架构后跟目标 ID 特征字符串的形式，由冒号分隔。每个目标 ID 特征是一个预定义的字符串，后跟加号或减号（例如gfx908:xnack+:sramecc-）。 此命令可以指定多次以启用多架构编译，例如 mcc -o axpy axpy.mu -lmusart --offload-arch=mp_21 --offload-arch=mp_22，此命令包括两对选项，指定架构 mp_21 和 mp_22，使编译器能够将 mp_21 和 mp_22 的二进制文件构建到一个胖二进制文件中。 如果没有添加任何卸载架构指定选项，则编译器将默认构建 mp_21 和 mp_22。

MUSA卸载目标ID：mp_10, mp_21, mp_22。

-mtgpu​

mcc通过后缀 .mu 识别 MUSA 源代码，需要通过此选项提示以控制编译将源代码编译到MTGPU目标。虽然mcc接受后缀为 .cu 的源文件，在这种情况下mcc需要选项 -mtgpu 来识别后缀为 .cu 的源文件作为MUSA语言源文件。

-cuda_wrapper​

mcc 可以在 CUDA 包装器模式下工作，支持开发人员使用 CUDA API 而不是 MUSA 的。

更多编译器选项可以参考 Clang 14.0.0 文档。

前端选项​

-Wno-shared-var-init​

MCC 支持使用 “-Wno-shared-var-init” 来忽略由共享内存初始化引起的编译错误，用户需要确保初始化操作是安全的。

后端选项​

每个后端选项都以 “-mllvm” 开头，表示以下选项是为后端准备的。

-mtgpu-maxregcnt= value​

为 MTGPU 目标设置最大临时寄存器计数。

-mtgpu-if-convert​

此选项将启用优化：“if” 块将被转换为小块大小的谓词执行。

-mtgpu-internalize-symbols= value​

此选项默认为 true，将尝试内联包括全局变量和函数在内的所有符号。

-mtgpu-enable-const-calc​

此选项默认为禁用，并启用常量计算程序生成优化。

-mtgpu-load-store-opt​

此选项默认为禁用，并将合并共享内存加载存储指令。

-enable-ldma-index​

此选项默认为禁用，并将启用 ldma 指令使用非零 imm-offset 索引进行指令选择阶段。

-mtgpu-prera-internal-opt​

此选项默认为禁用，并将启用优化，使用内部寄存器减少寄存器分配前临时寄存器的寄存器压力。此选项不能与 - mtgpu-postra-internal-opt 同时启用。

-mtgpu-postra-internal-opt​

此选项默认为禁用，并将启用优化，使用内部寄存器减少寄存器分配后临时寄存器的寄存器压力。此选项不能与 - mtgpu-prera-internal-opt 同时启用。

-internal-opt-avoid-use-nearest​

此选项默认为禁用，并将扩大内部寄存器优化时的优化间隔。通常与上述两个选项一起使用。

-mtgpu-load-cluster-mutation​

此选项默认为禁用，并将启用“加载集群变异”在寄存器分配后的指令顺序上。

-mtgpu-memory-sched-mutation​

此选项默认为禁用，并将启用“内存特定调度变异”在寄存器分配后的指令顺序上。

-mtgpu-tiny-offset-hint​

此选项默认为禁用，并将提示编译器代码不会使用大偏移量进行内存访问，这将超过 2^32 以获得更多的优化机会。

特殊控制选项​

-mathx-disable-fp64​

SUDI上没有硬件支持，所有 FP64 操作都由 libdevice 实现，这将扩大代码大小并导致编译时间过长。此选项将禁用任何 FP64 操作，并加速 SUDI（mp_10）的编译。

示例​

一个MUSA源代码文件：axpy.mu，使用以下命令行编译为可执行文件：

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

其中：

  • "axpy.mu"：是输入源文件。
  • "-lmusart"：链接 MUSA 运行时库，libmusart.so。
  • "-L/usr/local/musa/lib"：提示编译器在路径 "/usr/local/musa/lib" 下查找库。
  • "-O2"：表示优化级别。目前 mcc 仅支持 O2 及以上优化级别。并且 这是 mcc 在编译设备端代码时的默认优化级别。
  • "-o axpy"：指定输出文件名，“axpy”，以及在执行命令的位置下的输出位置。

如果源文件不以 .mu 结尾，则需要一对选项：-x musa，命令如下：

mcc -x musa axpy.cu -lmusart -L/usr/local/musa/lib -O2 -o axpy

如果源文件中的代码使用了 CUDA 的 API，如 cudaMemcpy，cudaMemset 等，请添加额外的选项 -cuda_wrapper 并链接 libcuda2musa.so 而不是 libmusart.so，命令如下：

mcc axpy.cu -mtgpu -cuda_wrapper -lcuda2musa -L/usr/local/musa/lib -O2 -o axpy

扩展属性

内核函数属性​

mtgpu_num_usreg​

MCC 支持为 MTGPU 目标使用 __attribute__((mtgpu_num_usreg(num_usreg))) 属性。 这些属性可以附加到内核函数定义上，是一个优化提示。

如果指定了这些属性，那么 MTGPU 目标后端将尝试将使用的 TEMP 寄存器和/或 ATTR 寄存器的数量限制为指定的值。使用的 TEMP 寄存器和/或 ATTR 寄存器的数量可能进一步向上取整以满足子目标的分配要求或约束。将 0 作为 num_usreg 意味着默认行为（无限制）。

用法：

__attribute__((mtgpu_num_usreg(256))) void KernelName(...) {...}

mtgpu_unroll_threshold​

MCC 支持为 MTGPU 目标使用 __attribute__((mtgpu_unroll_threshold(unroll_threshold))) 属性。这些属性可以附加到内核函数定义上，是一个优化提示。

用法：

__attribute__((mtgpu_unroll_threshold(400))) void KernelName(...) {...}

mtgpu_tiny_offset​

MCC 支持为 MTGPU 目标使用 __attribute__((mtgpu_tiny_offset)) 属性。这个属性可以附加到内核函数定义上，以提示编译器该内核中的内存访问操作的偏移量不会超过 INT_MAX，以创造更多的优化机会。

用法：

__global__ __attribute__((mtgpu_tiny_offset)) void kernel_name(...) {...}

版权所有 2024 摩尔线程，本文档受国际版权法保护。

摩尔线程和摩尔线程徽标是摩尔线程智能科技 (北京)有限责任公司的注册商标。

免责声明

本文档提供有关摩尔线程产品的信息。本文档并未授权任何知识产权的许可，并未以明示或暗示，或以禁止反言或其他方式授予任何知识产权许可。