毕昇异构算子开发全流程

毕昇异构算子开发的整体工作流,供项目成员参考。

毕昇异构算子开发全流程

前期准备

源码准备

先将源码仓库fork到自己的仓库中,并将自己的仓库克隆到本地,这里的本地指服务器。

1
git clone https://gitee.com/xxx/itk.git # xxx替换成自己的gitee用户名

克隆好之后,在与源码同级的文件夹中创建两个文件夹,用来存放编译构建和安装生成的文件。

1
2
mkdir ITK-build
mkdir ITK-install

截止目前文件结构应如下所示。

1
2
3
4
.
|--itk
|--ITK-build
|--ITK-install

git准备

进入源码文件夹,并新建一个分支。以我的算子为例,创建一个名为vnl_matrix_update的分支。

1
2
cd itk
git branch vnl_matrix_update

创建好后切换到该分支上。

1
git checkout vnl_matrix_update

可以看到成果切换分支的提示。

1
Switched to branch 'vnl_matrix_update'

然后提交分支。

1
2
git add . # 将所有修改添加至暂存区
git commit -m 'create branch' # 将暂存区的所有内容提交至本地仓库

如果在执行commit时出现如下提示:

1
2
3
4
5
6
7
8
9
10
11
*** Please tell me who you are.

Run

git config --global user.email "you@example.com"
git config --global user.name "Your Name"

to set your account's default identity.
Omit --global to set the identity only in this repository.

fatal: empty ident name (for <bisheng_tester2@csobluex.(none)>) not allowed

则根据提示对git进行配置。

1
2
git config --global user.email "you@example.com" # 邮箱替换为自己的
git config --global user.name "Your Name" # 用户名也替换为自己的

然后推送到远程仓库。

1
git push origin vnl_matrix_update # 将vnl_matrix_update分支推送到远程仓库origin

执行push后按照提示输入Gitee的用户名和密码。

1
2
3
4
5
6
7
8
Username for 'https://gitee.com': xxx # xxx替换成自己的gitee用户名
Password for 'https://yichu12138@gitee.com': # 输入密码
Total 0 (delta 0), reused 0 (delta 0)
remote: Powered by GITEE.COM [GNK-6.4]
remote: Create a pull request for 'vnl_matrix_update' on Gitee by visiting:
remote: https://gitee.com/yichu12138/itk/pull/new/yichu12138:vnl_matrix_update...yichu12138:master
To https://gitee.com/yichu12138/itk.git
* [new branch] vnl_matrix_update -> vnl_matrix_update

看到上述提示后,则表明分支已经成果提交到了自己的远程仓库中。

首次编译构建

按照之前的文档ITK线上环境编译与VSCode远程环境接入,将未改动过的源码编译构建并安装。

不建议先修改代码再编译构建!!!可能会遇到问题。

通用修改项

接下来先修改大家都需要进行修改的地方,这些修改主要是为了让ITK适配毕昇编译器。

顶层CMakeLists.txt

第46行修改前:

1
2
3
if(NOT CMAKE_CXX_STANDARD)
set(CMAKE_CXX_STANDARD 11) # Supported values are ``11``, ``14``, and ``17``.
endif()

修改后:

1
2
3
if(NOT CMAKE_CXX_STANDARD)
set(CMAKE_CXX_STANDARD 17) # 修改C++标准为17,源码为11
endif()

第231行修改前:

1
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${ITK_REQUIRED_CXX_FLAGS}")

修改后:

1
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${ITK_REQUIRED_CXX_FLAGS} -lsycl")

其他修改项

其余的修改项要根据自己的算子来做相应的调整。

CMakeLists.txt

找到与算子同级的CMakeLists.txt,即管理所修改算子的CMakeLists.txt

vnl_matrix算子为例,vnl_matrix.h的路径为Modules/ThirdParty/VNL/src/vxl/core/vnl/vnl_matrix.h,则管理它的CMakeLists.txt路径为Modules/ThirdParty/VNL/src/vxl/core/vnl/CMakeLists.txt

当然也不能完全靠路径来判断,在找到的CMakeLists.txt中应该能找到算子的名称之类的信息,来确定这个CMakeLists.txt到底是不是管理该算子。

例如vnl_matrix对应的CMakeLists.txt中就包含如下信息,从而可以确定。

找到对应的CMakeLists.txt后,在末尾添加一行代码。

1
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl -fsycl-targets=ascend_910-cce")

算子修改

找到要修改的算子,并找到算子的实现。

我要修改的update()算子位于Modules/ThirdParty/VNL/src/vxl/core/vnl/vnl_matrix.h:421,相应的实现位于Modules/ThirdParty/VNL/src/vxl/core/vnl/vnl_matrix.hxx:634

现在vnl_matrix.hxx中引入以下两个头文件。

1
2
#include <sycl/sycl.hpp>
#include <bisheng/bisheng.hpp>

然后对算子代码进行修改。修改前如下所示:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
template <class T>
vnl_matrix<T>& vnl_matrix<T>::update (vnl_matrix<T> const& m,
unsigned top, unsigned left)
{
unsigned int bottom = top + m.num_rows;
unsigned int right = left + m.num_cols;
#ifndef NDEBUG
if (this->num_rows < bottom || this->num_cols < right)
vnl_error_matrix_dimension ("update",
bottom, right, m.num_rows, m.num_cols);
#endif
for (unsigned int i = top; i < bottom; i++)
for (unsigned int j = left; j < right; j++)
this->data[i][j] = m.data[i-top][j-left];
return *this;
}

目前的代码暂未考虑模板类型和数据对齐的问题,只是为了展示整个工作流程。修改后如下所示:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
template <class T>
vnl_matrix<T>& vnl_matrix<T>::update (vnl_matrix<T> const& m,
unsigned top, unsigned left)
{
unsigned int bottom = top + m.num_rows;
unsigned int right = left + m.num_cols;
#ifndef NDEBUG
if (this->num_rows < bottom || this->num_cols < right)
vnl_error_matrix_dimension ("update",
bottom, right, m.num_rows, m.num_cols);
#endif

const auto thisM = this->num_rows;
const auto thisN = this->num_cols;
const auto mM = m.num_rows;
const auto mN = m.num_cols;
sycl::queue Q(sycl::ascend_selector{});

auto * devThis = malloc_device<int>(thisM * thisN, Q);
auto * devM = malloc_device<int>(mM * mN, Q);

Q.memcpy(devThis, *this->data, thisM * thisN * sizeof(int));
Q.memcpy(devM, *m.data, mM * mN * sizeof(int));

Q.launch<class Test>(mM, [=](sycl::group<1> group) {
__local int UBA[mN];
size_t groupId = group.get_id();
sycl::dmi::memcpy_blocks(UBA, &devM[groupId * mN], mN * sizeof(int) / 32);
sycl::dmi::memcpy_blocks(&devThis[top * thisN + groupId * thisN + left], UBA, mN * sizeof(int) / 32);
});

Q.memcpy(*this->data, devThis, thisM * thisN * sizeof(int));

Q.wait();

return *this;
}

具体的业务逻辑要根据自己算子的功能来编写。

修改后编译构建

算子修改完成后,进入ITK-build文件夹。由于刚才修改了CMakeLists.txt,故需执行一次cmake来刷新Makefile

1
cmake ../itk/

然后使用make编译构建并安装。

1
make -j24 && make install

安装好后就可以测试了,测试通过就可以认为完成了代码的修改工作。

提交成果

最后将修改后的成果提交至仓库中。

1
2
git add .
git commit -m 'operator vnl_matrix::update() heterogeneous code'

执行commit后会提示本次提交的代码与上次提交代码之间的差异。

1
2
[vnl_matrix_update a158479] operator vnl_matrix::update() heterogeneous code
3 files changed, 33 insertions(+), 5 deletions(-)

最后推送到远程仓库。

1
git push origin vnl_matrix_update

出现推送成果的提示即可。

1
2
3
4
5
6
7
8
Counting objects: 23, done.
Delta compression using up to 192 threads.
Compressing objects: 100% (12/12), done.
Writing objects: 100% (12/12), 1.45 KiB | 0 bytes/s, done.
Total 12 (delta 10), reused 0 (delta 0)
remote: Powered by GITEE.COM [GNK-6.4]
To https://gitee.com/yichu12138/itk.git
2fb588c..a158479 vnl_matrix_update -> vnl_matrix_update

发起Pull Request

登录Gitee到自己的仓库中,点击+ Pull Request发起PR。

输入标题即可提交,当然最好在说明中简略说明本次PR的目的和所修改的代码。

毕昇异构算子开发全流程

https://deleter-d.github.io/posts/35929/

作者

亦初

发布于

2023-05-20

更新于

2024-06-19

许可协议

评论

:D 一言句子获取中...

加载中,最新评论有1分钟缓存...