跳到主要内容

使用musify对代码进行平台迁移

· 阅读需 9 分钟

1. 引言

近些年,NVIDIA及其CUDA生态占据着GPGUP的领先地位,国产厂商也在奋力追赶。基于现实的考量,追赶者们需要尊重业界现存大量CUDA代码的事实,因此我们推出了musify这一工具,旨在尽量将原有的CUDA代码便捷无痛地转化,迁移至MUSA平台上编译运行,从而达到间接兼容的效果。

2. 如何使用musify

2.1. 安装依赖

musify使用python编写,要运行musify首先需要系统中安装有python3,然后假设用户使用ubuntu为例,执行以下命令

# 使得python默认指向python3
sudo apt install python-is-python3 -y
# 安装python依赖管理工具
sudo apt install pip -y
# 如果最后一行依赖库安装存在网络问题可以按需配置更新源,取消下一行的注释是一种可行的选择
# pip config set global.index-url https://pypi.mirrors.ustc.edu.cn/simple/
# 安装musify的依赖库
pip install ahocorapy

2.2. 帮助信息

和一般的命令行程序一样,musify通过-h/--help提供了帮助信息,有经验的用户可以直接参考帮助信息进行使用

$ musify-text -h
usage: musify-text [-h] [-t | -c | -i] [-d {c2m,m2c}] [-m [MAPPING ...]] [--clear-mapping] [-l {DEBUG,INFO,WARNING}] [srcs ...]

positional arguments:
srcs source files to be transformed

options:
-h, --help show this help message and exit
-t, --terminal print code to stdout
-c, --create write code to newly created file, default action
-i, --inplace modify code inplace
-d {c2m,m2c}, --direction {c2m,m2c}
convert direction
-m [MAPPING ...], --mapping [MAPPING ...]
api mapping
--clear-mapping clear default and previous mapping
-l {DEBUG,INFO,WARNING}, --log-level {DEBUG,INFO,WARNING}
lowest log level to display

如上信息,本工具的使用方式是先设置一系列的选项,最后跟着需要转化的代码文件列表

以下是对选项的详细解释:

-t, -c, -i是输出方式,使用时三选其一;分别为输出到屏幕,另外新建文件,直接修改原来的代码文件;默认行为是新建文件输出(但是如果代码有版本管理等保护,使用-i/--inplace直接原地修改比较方便)
-d设置方向c2m是cuda至musa,m2c是musa至cuda;默认为cuda至musa
-m指定替换使用的映射表,是映射表是json文件,内容为单层无嵌套的json object,其中每个name,value对分别表示相应的cuda和musa命名;一般不用指定使用默认即可,如果有特殊需求可自己生成映射表并指定,如果有多个映射表,使用如下格式-m a.json m b.json,即每个映射表前都有一个前置的-m
-m会添加新的映射表,不会覆盖原有的默认映射表,如果不想要默认映射表,需要使用--clear-mapping,然后在这个选项右边使用-m指定自己想要的映射表
-l指定日志等级,高于设置等级的日志都会被打印出来,默认为INFO
前面的选项顺序可以随意调整,如果使用默认值也可以不写,但是需要转化的代码文件必须跟在最后;同时为了防止存在以-开头的文件路径被识别为选项,在所有选项之后,文件路径之前,加上--

2.3. 使用示例

经过了如上的解释后,我们可以来到一个最简单的例子,使用默认映射表,默认转化方向(cuda至musa),只有输出方式改为了原地修改

如下命令修改了当前目录下的a.cpp和b.cpp两个文件

musify-text --inplace -- a.cpp b.cpp

但是一般而言,我们并不想手动列举需要转化的文件,一般一个项目有大量的代码文件需要转化,此时我们可以借助shell的``语法和一些第三方的文件搜索工具,比如

# 递归查找目录${DIR}下所有后缀名为cu,cuh,cpp或h的文件并转化
musify-text --inplace -- `find ${DIR} -name '*.cu' -name '*.cuh' -name '*.cpp' -name '*.h'`

同时我们可以推荐一些更加现代化的工具

# 安装此处使用的工具
sudo apt install ripgrep -y
# 如果没有自己编写kernel,cuda程序可以是纯C/C++代码,ripgrep对于cpp的后缀名自带预设,-tcpp就可以调用,rg --files则递归搜索指定目录下所有符合要求的文件
musify-text --inplace -- `rg --files -tcpp ${DIR}`
# ripgrep没有内置cuda的后缀模式,但是我们可以通过--type-add指定,然后再当场调用
musify-text --inplace -- `rg --files --type-add 'cuda:*.cu' --type-add 'cuda:*.cuh' -tcuda -tcpp ${DIR}`
# 如果嫌上面两种复杂,也可以使用-g选项直接指定后缀名,和前面的find作用基本一致
musify-text --inplace -- `rg --files -g '*.cu' -g '*.cuh' -g '*.cpp' -g '*.h' ${DIR}`

2.4. 排除标记

可能有不少读者注意到了musify的命令名称为musify-text,其实这是因为它的算法是单纯的文本匹配。

这种实现方案的原因主要是发现语法分析会引入过重的依赖(尤其对C++这种上下文相关文法),给用户的使用带来不便,纯粹文本匹配的工具更加小巧灵活。

当然这样带来了一定的转化可能不够智能的问题,于是我们加入了类似于lcov行覆盖率工具的排除标记功能,可以标记一定范围的代码不受转化影响,保持原样。

使用排除标记需要一定程度上修改原有代码,加入排除标记。

MUSIFY_EXCL_LINE
包含本标记的行会被排除
MUSIFY_EXCL_START
从包含本标记的行开始,直到MUSIFY_EXCL_STOP之间的行会被排除;当前行也会被排除
MUSIFY_EXCL_STOP
结束由MUSIFY_EXCL_START开启的排除;当前行不会被排除

其中START和STOP成对使用

// MUSIFY_EXCL_START
char *str_array[] = {
"cuInit",
"cuMalloc"
};
// MUSIFY_EXCL_STOP

LINE单独使用

char str[] = "cuInit"; // MUSIFY_EXCL_LINE

3. 总结与展望

本文介绍了musify的设计意图,使用方法和当前的缺陷。可能有不少读者也意识到如果没有其他不同版本需要加以区分,单纯因为使用了文本匹配就将工具叫做musify-text是不充分的;事实上,介于语法分析过重,文本匹配不够智能,目前存在一个后续计划是引入词法分析进行一定的代码分析识别,准备使用musify-lexer作为命令名称。

使用cmake构建MUSA工程

· 阅读需 21 分钟

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