使用cmake构建MUSA工程

摩尔线程

已于 2024-07-19 12:04:02 修改

阅读量558

点赞数 20

文章标签： c++ 开发语言

于 2024-05-26 09:49:05 首次发布

本文链接：https://blog.csdn.net/weixin_55064320/article/details/139207738

版权

更多内容请访问摩尔线程博客中心

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.txt内容如下：

#### 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编译选项
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前缀。

摩尔线程

关注

20
点赞
踩
20

收藏

觉得还不错? 一键收藏
0
评论
使用cmake构建MUSA工程

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

扫一扫