毕昇异构算子开发的整体工作流,供项目成员参考。
毕昇异构算子开发全流程
前期准备
源码准备
先将源码仓库 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 ) endif ()
修改后:
1 2 3 if (NOT CMAKE_CXX_STANDARD) set (CMAKE_CXX_STANDARD 17 ) 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
。
然后使用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的目的和所修改的代码。