CANN/catlass aclnn接口算子接入示例

basic_matmul_aclnn example【免费下载链接】catlass本项目是CANN的算子模板库提供NPU上高性能矩阵乘及其相关融合类算子模板样例。项目地址: https://gitcode.com/cann/catlassaclnn接口是CANN软件栈一直沿用的接口msOpGen工具是CANN提供可以生成该接口工程框架的工具便于用户编写一个具有aclnn接口的算子并使能CANN软件栈上的各种功能。该样例提供CATLASS算子模板接入msOpGen工程的示例代码与注意事项并提供CATLASS example风格的调用示例。下面以basic_matmul接入为例进行示例利用msOpGen工具接入该算子模板。1. 创建算子工程参考创建算子工程链接编写一个算子原型的json文件并生成对应工程。编写json相关示例代码catlass_basic_matmul.json生成工程执行下列脚本调用msOpGen生成算子工程。msopgen gen -i catlass_basic_matmul.json -c ai_core-soc_version -lan cpp -out catlass_basic_matmul其中soc_version可通过npu-smi info查看形如Ascendxxxyyy。需保证输入的json配置文件上例中的catlass_basic_matmul.json具有644的权限需保证输出的文件路径上例中的catlass_basic_matmul具有755的权限2. 编写Host代码参考Host侧Tiling实现-基本流程实现TilingFunc。若需要使能算子入图请参考算子入图GE图开发实现InferShape和InferDataType。相关示例代码 op_host/catlass_basic_matmul.cpp op_host/catlass_basic_matmul_tiling.h3. 编写Device代码参考Kernel侧算子实现实现kernel代码。相关示例代码 op_kernel/catlass_basic_matmul.cpp注意事项我们需要增加编译选项来引入CATLASS的头文件。在op_kernel/CMakeLists.txt中增加包含路径和架构宏选项。添加包含路径选项-I${CATLASS_INCLUDE_PATH}。其中${CATLASS_INCLUDE_PATH}是CATLASS代码仓下的include文件夹的路径需根据环境实际情况进行配置。添加架构宏选项-DCATLASS_ARCH${ARCH}。其中${ARCH}是对应架构的编号。根据CANN版本的不同默认写法有所不同CANN版本9.0.0.beta2# ... npu_op_kernel_options(ascendc_kernels ALL OPTIONS -I${CATLASS_INCLUDE_PATH}) # ...CANN版本9.0.0.beta2# set custom compile options if (${CMAKE_BUILD_TYPE}x STREQUAL Debugx) add_ops_compile_options(ALL OPTIONS -g -O0) endif() add_ops_compile_options(ALL OPTIONS -I${CATLASS_INCLUDE_PATH}) add_kernels_compile()msOpGen工程的分离编译模式不支持直接将结构体如Catlass::GemmCoord作为kernel的参数传入。当需要使用结构体时需要通过tiling地址传递成员数据然后在kernel侧重新构造。// 正确 extern C __global__ __aicore__ void catlass_basic_matmul(GM_ADDR self, GM_ADDR mat2, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); Catlass::GemmCoord problemShape{tiling_data.m, tiling_data.n, tiling_data.k}; // ... } // 暂不支持 extern C __global__ __aicore__ void catlass_basic_matmul(GM_ADDR self, GM_ADDR mat2, GM_ADDR out, GM_ADDR workspace, Catlass::GemmCoord problemShape) { // ... }4. 编译、部署参考算子工程编译、算子包部署进行编译、部署并设定环境变量。一般来说调用者需要添加头文件aclnn_catlass_basic_matmul.h并链接libcust_opapi.so。在不修改工程参数的情况下这两个文件的位置如下$ASCEND_HOME_PATH/opp/vendors/customize/op_api/include/aclnn_catlass_basic_matmul.h $ASCEND_HOME_PATH/opp/vendors/customize/op_api/lib/libcust_opapi.so这可作为Makefile/CMakeLists.txt的编写参考。可在5. 调用中查看CMake编写示例。5. 调用参考接口简介尝试调用。可参考以下内容编写CMakeLists.txtproject(basic_matmul_aclnn) cmake_minimum_required(VERSION 3.22) set(CATLASS_REPO_DIR 修改为实际环境上的CATLASS仓库路径) add_executable(basic_matmul_aclnn basic_matmul_aclnn.cpp) target_include_directories(basic_matmul_aclnn PRIVATE ${CATLASS_REPO_DIR}/examples/common ${CATLASS_REPO_DIR}/include $ENV{ASCEND_HOME_PATH}/include $ENV{ASCEND_HOME_PATH}/include/aclnn $ENV{ASCEND_HOME_PATH}/include/experiment/runtime $ENV{ASCEND_HOME_PATH}/include/experiment/msprof # 自定义算子包头文件目录 $ENV{ASCEND_HOME_PATH}/opp/vendors/customize/op_api/include ) target_link_directories(basic_matmul_aclnn PRIVATE $ENV{ASCEND_HOME_PATH}/lib64 # 自定义算子包库文件目录 $ENV{ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib/ ) target_link_libraries(basic_matmul_aclnn PRIVATE ascendcl nnopbase # 自定义算子包库文件名称 cust_opapi )预置示例我们对以上操作过程进行了集成以便快速体验aclnn工程接口。编译指定用例bash scripts/build.sh basic_matmul_aclnn cd output/run chmod x ./custom_opp_*.run ./custom_opp_*.run export LD_LIBRARY_PATH$ASCEND_HOME_PATH/opp/vendors/customize/op_api/lib/:${LD_LIBRARY_PATH} cd output/bin # 可执行文件名 |矩阵m轴|n轴|k轴|Device ID # Device ID可选默认为0 ./basic_matmul_aclnn 256 512 1024 0执行结果如下说明精度比对成功。Compare success.注意事项本示例仅用于CATLASS算子接入msopgen的参考为保证代码简洁不进行泛化的支持如多个算子、多个平台等。目前仅提供basic_matmul算子接入示例。示例仅支持以下产品Atlas A2 训练系列产品 / Atlas A2 推理系列产品(2201架构)Atlas A3 训练系列产品 / Atlas A3 推理系列产品(2201架构)【免费下载链接】catlass本项目是CANN的算子模板库提供NPU上高性能矩阵乘及其相关融合类算子模板样例。项目地址: https://gitcode.com/cann/catlass创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考