使用 cmake 构建 MUSA 工程

2024-05-20

1. 引言

cmake 是一个构建软件工程的工具，可以在不同的平台上使用统一的脚本生成相应平台对应的 Makefile 或者 build.ninja 文件。自 cmake 版本 2.8 之后，cmake 官方新增了 CUDA 的 Modules，使得可以使用 find_package(CUDA) 来获取构建 CUDA 项目所需要用到的 cmake 宏，变量以及函数。从 cmake 版本 3.9 之后，官方将 CUDA 加入了 cmake 的内置语言，从而可以使用 project(PROJECT_NAME LANGUAGES CXX CUDA) 或者 enable_language(CUDA) 来让使用 CUDA 语言。

由于 MUSA 暂时未被 cmake 官方收录，构建 MUSA 项目的方式将使用 Modules 的方案，即在 cmake 中使用 find_package(MUSA)。本文将详细介绍使用 cmake 构建 MUSA 工程的整体流程以及一些细节。

2. 简单的 MUSA 程序

我们从最简单的 MUSA 程序说起，项目目录如下，仅有一个 mu 代码文件。

-- project
 |-- main.mu

其中代码文件 main.mu 内容如下:

/* main.mu */
#include <iostream>
 
__global__ void func_kernel(float *x, float *y, float a) {
  y[threadIdx.x] = a * x[threadIdx.x];
}
 
int main(int argc, char *argv[]) {
  const int kDataLen = 4;
 
  float a = 2.0f;
  float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
  float host_y[kDataLen];
 
  // Copy input data to device.
  float *device_x;
  float *device_y;
  musaMalloc(&device_x, kDataLen * sizeof(float));
  musaMalloc(&device_y, kDataLen * sizeof(float));
  musaMemcpy(device_x, host_x, kDataLen * sizeof(float),
             musaMemcpyHostToDevice);
 
  // Launch the kernel.
  func_kernel<<<1, kDataLen>>>(device_x, device_y, a);
 
  // Copy output data to host.
  musaDeviceSynchronize();
  musaMemcpy(host_y, device_y, kDataLen * sizeof(float),
             musaMemcpyDeviceToHost);
 
  // Print the results.
  std::cout << "func" << std::endl;
  for (int i = 0; i < kDataLen; ++i) {
    std::cout << "y[" << i << "] = " << host_y[i] << "\n";
  }
 
  musaFree(device_x);
  musaFree(device_y);
  return 0;
}

为了节省空间，下面展示的代码有大段重复部分被省略，省略部分可从上面复制，省略后如下展示：

/* main.mu */
#include <iostream>
 
__global__ void func_kernel(float *x, float *y, float a) {...}
 
int main(int argc, char *argv[]) {
  ...
  func_kernel<<<1, kDataLen>>>(device_x, device_y, a);
  ...
  return 0;
}

编译代码只需要执行简单编译命令即可，编译和执行结果如下：

$ mcc main.mu -lmusart -o main
$ ./main
func
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8

main.mu 文件算是最小的 MUSA 代码，里面主函数执行了 GPU 程序的典型步骤：

申请显存， 将数据从 host 传输到 GPU 上， 执行 device 函数进行计算， 从 GPU 将数据取回 host， 释放显存。 其中展示的 device 上的计算任务是简单的向量缩放，对输入数据的每个元素乘以一个常数。代码文件后缀名为 .mu，编译器 mcc 会识别这个后缀名并以此为依据认为代码文件中包含 device 代码的定义和调用，即 __global__ 前缀的函数定义，和主函数中三尖括号 <<<...>>> 标记的函数调用，这两个是 MUSA 代码的最主要的标志，只能使用 mcc 编译器编译。

倘若代码文件命名为 main.cpp，即后缀为 .cpp，那么用上面的命令编译将会报错。原因是 .cpp 后缀默认约定指示该代码文件是常规的 c++ 代码并不包含 device 函数，这将自动调用 host 端的编译器如 g++ 执行编译。于是 g++ 将无法识别 MUSA 代码的语法而报错。这个时候需要执行的编译命令是

$ mcc -x musa main.cpp -lmusart -o main

其中需要在代码文件 main.cpp 的前面添加编译参数 -x musa, 这个编译参数告诉 mcc，虽然这个文件后缀名是 .cpp 但是它里面的内容是包含 MUSA 代码的，需要用 mcc 来执行编译。

3. 稍微复杂的工程项目

一个实际的项目，会有明确的组织结构，一般不会将 device 代码和 host 端的代码混合在一个代码文件中，否则不利于项目的维护。我们考虑一个典型的 device 代码和 host 代码分离的项目，其精简版的目录结构如下：

-- project
 |-- include
   |-- device_func.h
 |-- device
   |-- device_func.mu
 |-- main.cpp

其中host端代码如下：

/* main.cpp */
#include "device_func.h"
int main(int argc, char *argv[]) {
  device_func();
  return 0;
}

host 端代码不包含任何 device 的代码，对 GPU 的使用是通过封装好的 host 端的函数调用完成的。这些封装的使用 GPU 进行计算的函数接口声明在头文件 device_func.h 中：

/* include/device_func.h */
#pragma once
void device_func();

而对 GPU 进行计算的函数实现则统一放在另外一个部分，这个例子中是放在 device 目录中。这里的 device_func.mu 的内容则和前面介绍的基本相同：

/* device/device_func.mu */
#include <iostream>
 
__global__ void func_kernel(float *x, float *y, float a) {...}
 
void device_func() {
  ...
  func_kernel<<<1, kDataLen>>>(device_x, device_y, a);
  ...
}

这样的工程目录，编译项目可以使用如下步骤：

$ mkdir build
$ mcc ./device/device_func.mu -fPIC -c -o ./build/device_func.o
$ c++ ./build/device_func.o -fPIC -shared -o ./build/libdevice.so
$ c++ main.cpp ./build/libdevice.so -I ./include -L /usr/local/musa/lib -lmusart -o ./build/main
$ ./build/main
func
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8

工程项目编译过程往往会产生许多过程文件，我们先创建build目录来存放编译过程和结果的输出。第二步使用 mcc 编译器，将 device 目录下的 .mu 代码文件编译，这个过程会编译代码文件里面的 device 端 MUSA 代码。第三步我们将这些编译好的 GPU 相关代码整理成库文件供后续使用，这里演示生成动态链接库，因此在这个步骤和上一个步骤要使用 -fPIC 参数指示编译时按照地址无关的方式处理。最后第四步，编译和链接 host 端的代码。上面的项目编译流程是规范的，是干净的。使用 GPU 的加速代码一定是包含 MUSA 代码的，因此将他们归集到一个部分，编译时用 mcc 进行编译，然后生成库文件，可以是静态库也可以是动态库。GPU 函数提供的接口声明在头文件中供 host 端代码使用，而 host 代码的编写则如往常一样，包含接口头文件，直接调用接口，在生成可执行文件的链接阶段链接上 GPU 函数的库文件即可。这种项目的结构，对于一个从原本纯 CPU 的程序进行 GPU 加速扩展，是非常自然的，GPU 加速库可以独立编写，客户端程序仅仅是将原本 CPU 函数的接口调用改成相同功能的 GPU 接口调用。

4. 使用 MUSA 模块构建含 GPU 代码的项目

从上面章节可以看到，对于一个实际的 c++ 项目工程，代码文件往往是多个的，一般完整的项目构建流程，对于每一个代码文件需要一条编译命令生成 .o 目标文件。对于阶段性的每一个库文件目标，需要一条链接命令执行。对于最终的每一个可执行文件的生成也需要一条链接命令执行。并且库文件的生成依赖 .o 文件，可执行文件又依赖库文件或者 .o 文件，这要求以上的编译和链接命令需要按照某种合适的顺序执行。对于大项目而言，直接用编译链接命令来构建项目是非常繁琐的。因此诞生了 make，ninja 等构建项目的工具，这些构建工具是通过描述目标的依赖关系以及生成目标的命令来编织流程的。由于不同平台的编译命令有差别，且不同的构建工具如 make 或 ninja 的实现也有差别，于是又诞生了 cmake 用以统一这些。使用 cmake 可以仅编写同一套 CMakeLists.txt，生成不同平台不同编译工具使用的编译脚本，如 make 使用的 Makefile 或者 ninja 使用的 build.ninja。

我们进一步增加项目的复杂度，目录结构：

-- project
 |-- include
   |-- device_module1.h
   |-- device_module2.h
 |-- device
   |-- include
     |-- kernel.muh
   |-- device_module1.mu
   |-- device_module2.cpp
   |-- CMakeLists.txt
 |-- main1.cpp
 |-- main2.cpp
 |-- CMakeLists.txt

其中host端代码文件内容： /* main1.cpp */

#include "device_module1.h"
int main(int argc, char *argv[]) {
  mod1_func1();
  return 0;
}

/* main2.cpp */
#include "device_module1.h"
#include "device_module2.h"
int main(int argc, char *argv[]) {
  mod1_func2();
  mod2_func3();
  return 0;
}

接口文件如下：

/* include/device_module1.h */
#pragma once
void mod1_func1();
void mod1_func2();

/* include/device_module2.h */
#pragma once
void mod2_func3();
device端的代码文件有：

/* device/device_module1.mu */
#include <iostream>
#include "kernel.muh"
void mod1_func1() {
  ...
  func_kernel<1><<<1, kDataLen>>>(device_x, device_y, a);
  ...
  std::cout << "mod1_func1" << std::endl;
  ...
}
void mod1_func2() {
  ...
  func_kernel<2><<<1, kDataLen>>>(device_x, device_y, a);
  ...
  std::cout << "mod1_func2" << std::endl;
  ...
}

/* device/device_module2.cpp */
#include <iostream>
#include "kernel.muh"
void mod2_func3() {
  ...
  func_kernel<3><<<1, kDataLen>>>(device_x, device_y, a);
  ...
  std::cout << "mod2_func3" << std::endl;
  ...
}

device代码使用到的通用模板kernel放在kernel.muh

/* device/include/kernel.muh */
#pragma once
template<int s>
__global__ void func_kernel(float *x, float *y, float a) {
  y[threadIdx.x] = s * a * x[threadIdx.x];
}

项目的结构关系是，device 端代码会提供三个函数接口，其中模块一实现2个函数，模块二实现了1个函数，这两个模块会打包成一个动态库 libdevice.so 供 host 端的程序使用，而 host 端的两个不同的主程序分别调用了这三个接口。我们接下来用 cmake 来完成这个项目的构建。

项目的根 CMakeLists.tx 内容如下：

#### CMakeLists.txt
cmake_minimum_required(VERSION 3.10)
project(demo LANGUAGES CXX)
#### 在子项目中编译libdevice库，供后面链接使用
add_subdirectory(device)
#### 编译程序1
add_executable(main1 main1.cpp)
target_include_directories(main1 PRIVATE
  ${CMAKE_CURRENT_SOURCE_DIR}/include
  )
target_link_libraries(main1 PRIVATE device)
#### 编译程序2
add_executable(main2 main2.cpp)
target_include_directories(main2 PRIVATE
  ${CMAKE_CURRENT_SOURCE_DIR}/include
  )
target_link_libraries(main2 PRIVATE device)

常规的手段编译两个主程序，注意要指定接口的头文件目录，以及需要指定需链接的 GPU 库，这个 GPU 库是在 device 子模块中编译生成的。

项目子目录 device 是一个子项目，里面同样包含一个 CMakeLists.txt 用于构建 GPU 库：

#### device/CMakeLists.txt
cmake_minimum_required(VERSION 3.10)
project(device LANGUAGES CXX)
#### 载入cmake的MUSA模块
list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)
find_package(MUSA REQUIRED)
#### 编译musa代码时打印信息
set(MUSA_VERBOSE_BUILD ON)
#### 添加额外的mcc编译选项, 下面例子指定编译成春晓架构的GPU代码，并将警告当成错误。
#### 注意，若要指定其他架构，将mp_21替换成相应架构号。
#### 例如MTT S50对应mp_10，MTT S80和MTT S3000对应mp_21，MTT S4000对应mp_22。
#### mcc 编译器的编译选项细节请参考编译器文档。
set(MUSA_MCC_FLAGS --offload-arch=mp_21 -Werror)
#### 添加头文件路径
musa_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include)
#### 让.cpp后缀的文件能被当成MUSA代码从而使用mcc编译
#### 若是mu或者cu后缀则无需指定，会自动被识别成MUSA代码
set_source_files_properties(device_module2.cpp
      PROPERTIES
      MUSA_SOURCE_PROPERTY_FORMAT OBJ
  )
#### 添加动态库libdevice.so
musa_add_library(device SHARED
  device_module1.mu
  device_module2.cpp
  )

一开始需要使用 find_package(MUSA REQUIRED)，来载入 MUSA 模块。这里有一点需要注意，由于该模块暂时未被 cmake 官方收录，仅安装在了 MUSA Toolkit 的安装目录中，因此需要在载入模块之前将模块的安装目录更新到 cmake 的 MODULE 搜索路径中: list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)。接下来就可以使用 musa_add_library 这个在 MUSA 模块定义的 cmake 函数宏来指定为这个项目添加一个库目标。里面输入的所有源文件都会为其挨个生成编译命令，最后生成链接命令来打包成库。若需要为编译时提供头文件路径，则使用 musa_include_directories 函数宏。需要注意的是，源文件列表中的代码文件，若后缀名是 .mu或者 .cu，则会自动被识别用 mcc 编译器编译。在这个例子中故意将 device_module2 这个代码文件的后缀名写成 .cpp，来模拟一个情况。正如第二章所说的，虽然代码文件后缀名是 .cpp 但是里面却含有 MUSA 代码，编译命令需要使用 mcc 编译器，并且加上 -x musa 编译参数。在 cmake 中，对于 .cpp 后缀文件同样按默认约定是当成常规的 c++ 代码文件的，默认使用 c++ 编译器。这里为了明确告知 cmake 这个文件包含 MUSA 代码，可以用 set_source_files_properties(device_module2.cpp PROPERTIES MUSA_SOURCE_PROPERTY_FORMAT OBJ) 来设置该代码文件的文件属性。这样 cmake 就会把这个代码文件等同于 .mu 后缀来处理。

MUSA的模块中还提供了 musa_add_executable 函数宏来直接生成可执行文件。例如，对于第二章中的简单项目而言，可以不生成中间的库文件，而是直接将所有 .o 目标文件链接成最后的可执行文件，这样使用 cmake 可以如下简单实现：

#### CMakeLists.txt
cmake_minimum_required(VERSION 3.10)
project(demo LANGUAGES CXX)
#### 载入cmake的MUSA模块
list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)
find_package(MUSA REQUIRED)
#### 编译主程序
musa_add_executable(main main.cpp ./device/device_func.mu)
target_include_directories(main PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)

在musa_add_executable 函数中的源文件，main.cpp 未作特殊处理，会被当成常规 c++ 文件被 c++ 编译器编译，而 device_func.mu 则会被用 mcc 编译器编译，最后两者被链接成可执行文件。

还有一些细节需要注意，musa_add_executable及musa_add_library 是自定义的函数宏并不是官方的 cmake 命令，它内部的实现使用了自定义命令来调用 mcc 编译器。于是为了传递信息给 mcc 编译器不能使用常规的 cmake 命令，比方设置头文件路径就不能直接使用 include_directories 或者 target_include_directories，这个只会传递给 c++ 编译器，若要传递给 mcc 编译器则使用 musa_include_directories 这个命令。设置编译选项，不能直接用 target_compile_options，而需要使用 set(MUSA_MCC_FLAGS --offload-arch=mp_21 -Werror)，通过 MUSA_MCC_FLAGS 这个变量传递给 mcc 编译器。倘若要显示编译过程信息，需要使用 set(MUSA_VERBOSE_BUILD ON) 这个开关来打开。

这个案例的编译命令和运行结果如下：

$ cmake -B build
$ cmake --build build
$ ./build/main1
mod1_func1
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8
$ ./build/main2
mod1_func2
y[0] = 4
y[1] = 8
y[2] = 12
y[3] = 16
mod2_func3
y[0] = 6
y[1] = 12
y[2] = 18
y[3] = 24

5. 使用MUSAToolkit模块构建项目

cmake 版本3.17之后，官方新增加了 CUDAToolkit 模块。其用途是，有些工程并没有包含任何 GPU 端的 kernel 代码，但是调用了 NVIDIA 官方提供的数学库或图像处理库等现成的 c 或 c++ 接口，从而使用 GPU 进行加速。这种情况下整个工程内全部代码都是 c 或 c++ 语言编写，无任何 CUDA 代码，故无需使用 nvcc 编译器编译工程。那么整个项目仅需要使用 host 端的 c 或 c++ 编译器编译，最后链接的时候把运行时库以及数学库等添加上即可。CUDAToolkit 模块则提供了所有使用 GPU 加速可能使用到的库目标。为了保证对 CUDA 使用的最佳兼容，MUSAToolkit 中也包含了 cmake 的 MUSAToolkit 模块。

例如如下的简单项目：

-- project
 |-- main.cpp
 |-- CMakeLists.txt

其中主函数里面是使用 mufft 数学库进行傅里叶变化计算：

/* main.cpp */
#include <iostream>
#include <mufft.h>
#include <musa_runtime.h>
int main() {
  const int Nx = 8;
  size_t complex_bytes = sizeof(float) * 2 * Nx;
  // create and initialize host data
  float *h_x = (float *)malloc(complex_bytes);
  for (size_t i = 0; i < Nx; i++) {
    h_x[2 * i] = i;
    h_x[2 * i + 1] = i;
  }
  // Create MUSA device object and copy data to device
  void *d_x;
  musaMalloc(&d_x, complex_bytes);
  musaMemcpy(d_x, h_x, complex_bytes, musaMemcpyHostToDevice);
  // Create the plan
  mufftHandle plan = NULL;
  mufftPlan1d(&plan, Nx, MUFFT_C2C, 1);
  // Execute plan:
  mufftExecC2C(plan, (mufftComplex *)d_x, (mufftComplex *)d_x, MUFFT_FORWARD);
  // copy back the result to host
  musaMemcpy(h_x, d_x, complex_bytes, musaMemcpyDeviceToHost);
  for (size_t i = 0; i < Nx; i++) {
    std::cout << "(" << h_x[2 * i] << ", " << h_x[2 * i + 1] << ")\n";
  }
  // release resource
  mufftDestroy(plan);
  musaFree(d_x);
  free(h_x);
  return 0;
}

该项目的 CMakeListst.txt 如下：

#### CMakeLists.txt
cmake_minimum_required(VERSION 3.10)
project(demo LANGUAGES CXX)
## 载入MUSAToolkit模块
list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)
find_package(MUSAToolkit REQUIRED)
## 添加可执行文件
add_executable(main main.cpp)
## 为目标链接运行时库及数学库
target_link_libraries(main PRIVATE
  MUSA::musart
  MUSA::mufft
  )

在一开始需要用 find_package(MUSAToolkit REQUIRED)，来载入 MUSAToolkit 模块。同样的由于该模块暂时未被 cmake 官方收录，仅安装在了 MUSA Toolkit 的安装目录中，因此需要在载入模块之前将模块的安装目录更新到 cmake 的 MODULE 搜索路径中: list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)。

模块载入之后，将提供若干库目标以及变量供使用，这个例子中用到了运行时库和傅里叶变换库，故给目标添加链接库 MUSA::musart和MUSA::mufft。需要指出的是，这里模块提供的目标 MUSA:: 已经包含了所需的头文件路径，会自动传递给要编译的目标，故无需再给编译目标添加 MUSA 相关的头文件目录。

6. 总结

MUSA 沿用了 Modules 的方式，也提供了相似的 cmake 模块供使用。保留了和 CUDA 几乎完全一致的使用方式，以达到用户尽可能方便地构建 MUSA 工程。这个兼容性也能带来快速迁移 CUDA 项目的便捷。在做项目迁移时，若项目使用 cmake 工具构建，则绝大多数情况下可以仅做文本替换，将 CMakelist.txt 中的 CUDA 替换成 MUSA，CU 前缀替换成 MU 前缀。