diff --git a/examples/index/arithprogression/README.md b/examples/index/arange/README.md similarity index 72% rename from examples/index/arithprogression/README.md rename to examples/index/arange/README.md index 61cdcd5fa58142f157c474b23f5df3e669e93a83..ec83a5e5ac2d0af736fa06c0f5b0f991bf0e969d 100644 --- a/examples/index/arithprogression/README.md +++ b/examples/index/arange/README.md @@ -2,13 +2,13 @@ ## 概述 -本样例介绍了调用ArithProgression高阶API实现arithprogression单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 +本样例介绍了调用Arange高阶API实现arange单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 -- 直调：使用核函数直调arithprogression自定义算子。 +- 直调：使用核函数直调arange自定义算子。 核函数的基础调用方式，开发者完成算子核函数的开发和Tiling实现后，即可通过AscendCL运行时接口，完成算子的调用。 -- 框架调用：使用框架调用arithprogression自定义算子。 +- 框架调用：使用框架调用arange自定义算子。 按照工程创建->算子实现->编译部署->算子调用的流程完成算子开发。整个过程都依赖于算子工程：基于工程代码框架完成算子核函数的开发和Tiling实现，通过工程编译脚本完成算子的编译部署，继而实现单算子调用或第三方框架中的算子调用。 @@ -17,7 +17,7 @@ | 调用方式 | 目录 | **描述** | | --------- | ------------------------------------------------------------ | ---------------------------------------------------------- | | 直调 | [kernel_launch_method_by_direct](./kernel_launch_method_by_direct) | host侧的核函数调用程序，包含CPU侧、NPU侧、仿真侧三种运行验证方法。 | -| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用arithprogression算子。 | +| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用arange算子。 | ## 样例支持的产品型号为： - Atlas A2训练系列产品/Atlas 800I A2推理产品 @@ -34,12 +34,12 @@ ## 算子描述 -arithprogression单算子，给定起始值，等差值和长度，返回一个等差数列。 +arange单算子，给定起始值，等差值和长度，返回一个等差数列。 -arithprogression算子规格： +arange算子规格： - + @@ -54,14 +54,14 @@ arithprogression算子规格： ## 算子实现介绍 -本样例中实现的是固定shape，输入firstGm[1]，diffGm[1]，输出outputGm[128]的arithprogression算子。 +本样例中实现的是固定shape，输入firstGm[1]，diffGm[1]，输出outputGm[128]的arange算子。 - kernel实现 - 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用ArithProgression高阶API接口完成arithprogression计算，得到最终结果，再搬出到外部存储上。 + 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Arange高阶API接口完成arange计算，得到最终结果，再搬出到外部存储上。 - arithprogression算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入搬运至LocalMemory，Compute任务负责根据输入信息执行arithprogression计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dst_global中。 + arange算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入搬运至LocalMemory，Compute任务负责根据输入信息执行arange计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dst_global中。 - tiling实现 - arithprogression算子的tiling实现流程如下：设置输出shape信息。 \ No newline at end of file + arange算子的tiling实现流程如下：设置输出shape信息。 \ No newline at end of file diff --git a/examples/index/arithprogression/host_tiling/arithprogression_custom_tiling.h b/examples/index/arange/host_tiling/arange_custom_tiling.h similarity index 56% rename from examples/index/arithprogression/host_tiling/arithprogression_custom_tiling.h rename to examples/index/arange/host_tiling/arange_custom_tiling.h index 1daa71c35dece4521b50b942b150077747639e6f..251a3d970bb67a1ecc0d941d8f24a91863eed26a 100644 --- a/examples/index/arithprogression/host_tiling/arithprogression_custom_tiling.h +++ b/examples/index/arange/host_tiling/arange_custom_tiling.h @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -14,14 +15,14 @@ #include "tiling/tiling_api.h" namespace optiling { -BEGIN_TILING_DATA_DEF(ArithprogressionCustomTilingData) +BEGIN_TILING_DATA_DEF(ArangeCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, outSize); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(ArithprogressionCustom, ArithprogressionCustomTilingData) +REGISTER_TILING_DATA_CLASS(ArangeCustom, ArangeCustomTilingData) } // namespace optiling -void ComputeTiling(const uint32_t outSize, optiling::ArithprogressionCustomTilingData &tiling){ +void ComputeTiling(const uint32_t outSize, optiling::ArangeCustomTilingData &tiling){ tiling.set_outSize(outSize); } diff --git a/examples/index/arithprogression/kernel_impl/arithprogression_kernel.h b/examples/index/arange/kernel_impl/arange_kernel.h similarity index 78% rename from examples/index/arithprogression/kernel_impl/arithprogression_kernel.h rename to examples/index/arange/kernel_impl/arange_kernel.h index 5df32c0bbbb41cb6b4d93b04b5bc4eb3dad41ffa..394abfa0f2e129622524e1d4d0d63411f7f93c06 100644 --- a/examples/index/arithprogression/kernel_impl/arithprogression_kernel.h +++ b/examples/index/arange/kernel_impl/arange_kernel.h @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -19,9 +20,9 @@ struct VecTiling{ }; template -class KernelArithProgression { +class KernelArange { public: - __aicore__ inline KernelArithProgression() + __aicore__ inline KernelArange() {} __aicore__ inline void Init(GM_ADDR dstGm, float firstValue, float diffValue, uint32_t count) { @@ -44,7 +45,7 @@ private: __aicore__ inline void Compute() { AscendC::LocalTensor dstLocal = outDst.AllocTensor(); - AscendC::ArithProgression(dstLocal, static_cast(firstValue_), static_cast(diffValue_), count_); + AscendC::Arange(dstLocal, static_cast(firstValue_), static_cast(diffValue_), count_); outDst.EnQue(dstLocal); } __aicore__ inline void CopyOut() diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_direct/CMakeLists.txt similarity index 80% rename from examples/index/arithprogression/kernel_launch_method_by_direct/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_direct/CMakeLists.txt index 77743ac5d9a64187d5ae894a9b4666c41d34445a..65f74305c96da4b1e68d06b60a14fe51506a9619 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/CMakeLists.txt +++ b/examples/index/arange/kernel_launch_method_by_direct/CMakeLists.txt @@ -18,7 +18,7 @@ if(CMAKE_INSTALL_PREFIX STREQUAL /usr/local) endif() file(GLOB KERNEL_FILES - ${CMAKE_CURRENT_SOURCE_DIR}/arithprogression_custom.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/arange_custom.cpp ) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") @@ -30,29 +30,29 @@ else() message("invalid RUN_MODE: ${RUN_MODE}") endif() -add_executable(arithprogression_direct_kernel_op +add_executable(arange_direct_kernel_op ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/arithprogression_custom_tiling.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/arange_custom_tiling.cpp ) -target_compile_options(arithprogression_direct_kernel_op PRIVATE +target_compile_options(arange_direct_kernel_op PRIVATE $:-g>> -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 ) -target_compile_definitions(arithprogression_direct_kernel_op PRIVATE +target_compile_definitions(arange_direct_kernel_op PRIVATE $<$>:CUSTOM_ASCEND310P> ) -target_include_directories(arithprogression_direct_kernel_op PRIVATE +target_include_directories(arange_direct_kernel_op PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} $:${ASCEND_CANN_PACKAGE_PATH}/include>> $:${ASCEND_CANN_PACKAGE_PATH}/runtime/include>> ) -target_link_libraries(arithprogression_direct_kernel_op PRIVATE +target_link_libraries(arange_direct_kernel_op PRIVATE $,$>:host_intf_pub>> $:tikicpulib::${SOC_VERSION}>> $:ascendcl>> @@ -66,7 +66,7 @@ target_link_libraries(arithprogression_direct_kernel_op PRIVATE graph_base ) -install(TARGETS arithprogression_direct_kernel_op +install(TARGETS arange_direct_kernel_op LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/README.md b/examples/index/arange/kernel_launch_method_by_direct/README.md similarity index 87% rename from examples/index/arithprogression/kernel_launch_method_by_direct/README.md rename to examples/index/arange/kernel_launch_method_by_direct/README.md index 3d1102c04b0c1f91758eb28c4e25621aa3b12047..2998b9424838c37f0f8014f29391325ae4515952 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/README.md +++ b/examples/index/arange/kernel_launch_method_by_direct/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于Kernel直调算子工程，介绍了调用ArithProgression高阶API实现arithprogression单算子，主要演示ArithProgression高阶API在Kernel直调工程中的调用。 +本样例基于Kernel直调算子工程，介绍了调用Arange高阶API实现arange单算子，主要演示Arange高阶API在Kernel直调工程中的调用。 ## 目录结构介绍 | 目录及文件 | 描述 | @@ -10,8 +10,8 @@ | [cmake](./cmake) | 编译工程文件 | | [scripts](./scripts) | 包含输入数据和真值数据生成脚本文件 | | main.cpp | 主函数，调用算子的应用程序，含CPU域及NPU域调用 | -| arithprogression_custom.cpp | 算子kernel实现 | -| arithprogression_custom_tiling.cpp | 算子tiling实现 | +| arange_custom.cpp | 算子kernel实现 | +| arange_custom_tiling.cpp | 算子tiling实现 | | CMakeLists.txt | 编译工程文件 | | run.sh | 编译执行脚本 | diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom.cpp b/examples/index/arange/kernel_launch_method_by_direct/arange_custom.cpp similarity index 77% rename from examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom.cpp rename to examples/index/arange/kernel_launch_method_by_direct/arange_custom.cpp index cb4aa938df814a818b7f5df1dcee50a244110d33..b90764c036af51a42794361d22052fe898fa33f9 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom.cpp +++ b/examples/index/arange/kernel_launch_method_by_direct/arange_custom.cpp @@ -1,14 +1,15 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "kernel_operator.h" -#include "../kernel_impl/arithprogression_kernel.h" +#include "../kernel_impl/arange_kernel.h" __aicore__ inline void CopyTiling(MyCustomKernel::VecTiling* tiling, GM_ADDR tilingGM) { @@ -27,7 +28,7 @@ extern "C" __global__ __aicore__ void arith_progression_custom(GM_ADDR firstGm, CopyTiling(&tilingData, tiling); auto first = reinterpret_cast<__gm__ float*>(firstGm); auto diff = reinterpret_cast<__gm__ float*>(diffGm); - MyCustomKernel::KernelArithProgression op; + MyCustomKernel::KernelArange op; op.Init(dstGm, *first, *diff, tilingData.outSize); op.Process(); } diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom_tiling.cpp b/examples/index/arange/kernel_launch_method_by_direct/arange_custom_tiling.cpp similarity index 54% rename from examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom_tiling.cpp rename to examples/index/arange/kernel_launch_method_by_direct/arange_custom_tiling.cpp index e16e6aebc469d775868f913164f6b3a3a2b5ff91..f319835b51444e363b22ba57ca5f1885010ba776 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/arithprogression_custom_tiling.cpp +++ b/examples/index/arange/kernel_launch_method_by_direct/arange_custom_tiling.cpp @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -14,17 +15,17 @@ #include #include #include "tiling/tiling_api.h" -#include "../host_tiling/arithprogression_custom_tiling.h" +#include "../host_tiling/arange_custom_tiling.h" -uint8_t *GetTilingBuf(optiling::ArithprogressionCustomTilingData *tilingData) { - uint32_t tilingSize = sizeof(optiling::ArithprogressionCustomTilingData); +uint8_t *GetTilingBuf(optiling::ArangeCustomTilingData *tilingData) { + uint32_t tilingSize = sizeof(optiling::ArangeCustomTilingData); uint8_t *buf = (uint8_t *)malloc(tilingSize); tilingData->SaveToBuffer(buf, tilingSize); return buf; } uint8_t* GenerateTiling(uint32_t outSize){ - optiling::ArithprogressionCustomTilingData tiling; + optiling::ArangeCustomTilingData tiling; ComputeTiling(outSize, tiling); return GetTilingBuf(&tiling); } \ No newline at end of file diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/cmake/cpu_lib.cmake b/examples/index/arange/kernel_launch_method_by_direct/cmake/cpu_lib.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_direct/cmake/cpu_lib.cmake rename to examples/index/arange/kernel_launch_method_by_direct/cmake/cpu_lib.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/cmake/npu_lib.cmake b/examples/index/arange/kernel_launch_method_by_direct/cmake/npu_lib.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_direct/cmake/npu_lib.cmake rename to examples/index/arange/kernel_launch_method_by_direct/cmake/npu_lib.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/main.cpp b/examples/index/arange/kernel_launch_method_by_direct/main.cpp similarity index 94% rename from examples/index/arithprogression/kernel_launch_method_by_direct/main.cpp rename to examples/index/arange/kernel_launch_method_by_direct/main.cpp index b9a1645d0ea8c35fbd3ef162c348b8d3a9a63315..ea4bcb412c722ee19effc473f86cfc2a716f53c7 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/main.cpp +++ b/examples/index/arange/kernel_launch_method_by_direct/main.cpp @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/run.sh b/examples/index/arange/kernel_launch_method_by_direct/run.sh similarity index 45% rename from examples/index/arithprogression/kernel_launch_method_by_direct/run.sh rename to examples/index/arange/kernel_launch_method_by_direct/run.sh index c49744fd3a5a83ab326655ffd23c09ef7f8edcb4..e94be31ac00ea489a29a4bb72303a9b3886aad9d 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/run.sh +++ b/examples/index/arange/kernel_launch_method_by_direct/run.sh @@ -26,6 +26,21 @@ rm -rf build mkdir build cd build +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== +if [ "${RUN_MODE}" = "sim" ]; then + export LD_LIBRARY_PATH=$(echo $LD_LIBRARY_PATH | sed 's/\/.*\/runtime\/lib64://g') + export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/runtime/lib64/stub:$LD_LIBRARY_PATH +fi + +source $ASCEND_HOME_DIR/bin/setenv.bash export LD_LIBRARY_PATH=${ASCEND_HOME_DIR}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH set -euo pipefail @@ -33,11 +48,11 @@ cmake -DRUN_MODE=${RUN_MODE} -DSOC_VERSION=${SOC_VERSION} -DASCEND_CANN_PACKAG make -j16 if [ "${RUN_MODE}" = "npu" ]; then - ./arithprogression_direct_kernel_op + ./arange_direct_kernel_op elif [ "${RUN_MODE}" = "sim" ]; then export ASCEND_TOOLKIT_HOME=${ASCEND_HOME_DIR} export ASCEND_HOME_PATH=${ASCEND_HOME_DIR} - msprof op simulator --application=./arithprogression_direct_kernel_op + msprof op simulator --application=./arange_direct_kernel_op elif [ "${RUN_MODE}" = "cpu" ]; then - ./arithprogression_direct_kernel_op + ./arange_direct_kernel_op fi diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/gen_data.py b/examples/index/arange/kernel_launch_method_by_direct/scripts/gen_data.py similarity index 70% rename from examples/index/arithprogression/kernel_launch_method_by_framework/scripts/gen_data.py rename to examples/index/arange/kernel_launch_method_by_direct/scripts/gen_data.py index c10122ce59c08d6ee43132518d0e03cfb401af3d..11c542cc45ca8ae51e1d92ac62d0d0da07eeb138 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/gen_data.py +++ b/examples/index/arange/kernel_launch_method_by_direct/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 -# Copyright (c) 2024 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import numpy as np import os diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_framework/CMakeLists.txt similarity index 81% rename from examples/index/arithprogression/kernel_launch_method_by_framework/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_framework/CMakeLists.txt index 584132d80993d309434fb1303de83910a1989aba..83a789d3366aa7794f09d01ba16d73c0aeac7f69 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/CMakeLists.txt +++ b/examples/index/arange/kernel_launch_method_by_framework/CMakeLists.txt @@ -33,7 +33,14 @@ if(ENABLE_TEST AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/testcases) add_subdirectory(testcases) endif() -# modify vendor_name in install.sh and upgrade.sh +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/scripts/install.sh ${CMAKE_BINARY_DIR}/scripts/upgrade.sh COMMAND mkdir -p ${CMAKE_BINARY_DIR}/scripts COMMAND cp -r ${CMAKE_SOURCE_DIR}/scripts/* ${CMAKE_BINARY_DIR}/scripts/ diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/CMakePresets.json b/examples/index/arange/kernel_launch_method_by_framework/CMakePresets.json similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/CMakePresets.json rename to examples/index/arange/kernel_launch_method_by_framework/CMakePresets.json diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/README.md b/examples/index/arange/kernel_launch_method_by_framework/README.md similarity index 92% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/README.md rename to examples/index/arange/kernel_launch_method_by_framework/README.md index 3f12aeeecb40af9c89d74e79e71b0496bc125282..95930a17808fbb06d73debcd5f189e3e27495ff5 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/README.md +++ b/examples/index/arange/kernel_launch_method_by_framework/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于自定义算子工程，介绍了调用SelectWithBytesMask高阶API实现SelectWithBytesMaskCustom单算子，主要演示SelectWithBytesMask高阶API在自定义算子工程中的调用。 +本样例基于自定义算子工程，介绍了调用Arange高阶API实现arange单算子，主要演示Arange高阶API在自定义算子工程中的调用。 ## 目录结构 | 目录 | 描述 | @@ -45,7 +45,7 @@ …… "ASCEND_CANN_PACKAGE_PATH": { "type": "PATH", - //请替换为CANN开发套件包安装后文件存储路径。例如：/home/HwHiAiUser/Ascend/ascend-toolkit/latest + // 请替换为CANN开发套件包安装后文件存储路径。例如：/home/HwHiAiUser/Ascend/ascend-toolkit/latest "value": "~/Ascend/ascend-toolkit/latest" }, …… @@ -75,7 +75,7 @@ cd build_out ### 5.执行样例 在build_out目录下执行如下命令 ``` -./select_with_bytes_mask_custom_npu +./arange_custom_npu ``` ### 6.NPU仿真模式运行（可选） 若要执行NPU仿真，在build_out目录下执行如下命令： diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/build.sh b/examples/index/arange/kernel_launch_method_by_framework/build.sh similarity index 95% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/build.sh rename to examples/index/arange/kernel_launch_method_by_framework/build.sh index 0813af07bbc6da272e38b7019e81b538bbbf8b30..229910e0133e72beadd6055c3ebb269442a6ea40 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/build.sh +++ b/examples/index/arange/kernel_launch_method_by_framework/build.sh @@ -3,7 +3,6 @@ script_path=$(realpath $(dirname $0)) source $ASCEND_HOME_DIR/bin/setenv.bash cp -rf ../host_tiling/* op_host/ -rm -rf ./cmake/util ln -s $ASCEND_HOME_DIR/tools/op_project_templates/ascendc/customize/cmake/util/ ./cmake/util mkdir -p build_out rm -rf build_out/* @@ -70,8 +69,3 @@ else fi -# for debug -# cd build_out -# make -# cpack -# verbose append -v \ No newline at end of file diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/cmake/config.cmake b/examples/index/arange/kernel_launch_method_by_framework/cmake/config.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/cmake/config.cmake rename to examples/index/arange/kernel_launch_method_by_framework/cmake/config.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/cmake/func.cmake b/examples/index/arange/kernel_launch_method_by_framework/cmake/func.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/cmake/func.cmake rename to examples/index/arange/kernel_launch_method_by_framework/cmake/func.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/cmake/intf.cmake b/examples/index/arange/kernel_launch_method_by_framework/cmake/intf.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/cmake/intf.cmake rename to examples/index/arange/kernel_launch_method_by_framework/cmake/intf.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/cmake/makeself.cmake b/examples/index/arange/kernel_launch_method_by_framework/cmake/makeself.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/cmake/makeself.cmake rename to examples/index/arange/kernel_launch_method_by_framework/cmake/makeself.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/cmake/util b/examples/index/arange/kernel_launch_method_by_framework/cmake/util similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/cmake/util rename to examples/index/arange/kernel_launch_method_by_framework/cmake/util diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/op_host/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_framework/op_host/CMakeLists.txt similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/op_host/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_framework/op_host/CMakeLists.txt diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom.cpp b/examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom.cpp similarity index 80% rename from examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom.cpp rename to examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom.cpp index 2d01f4dc704b79ceb9086c55146cb6cb4cf849e5..3f87142a1cbeb72c849eae9ec888e0b330643185 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom.cpp +++ b/examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom.cpp @@ -1,21 +1,22 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "arithprogression_custom_tiling.h" +#include "arange_custom_tiling.h" #include "register/op_def_registry.h" namespace optiling { constexpr uint32_t BLOCK_DIM = 1; static ge::graphStatus TilingFunc(gert::TilingContext* context) { - ArithprogressionCustomTilingData tiling; + ArangeCustomTilingData tiling; const gert::StorageShape* x_shape = context->GetOutputShape(0); const gert::Shape shape = x_shape->GetStorageShape(); @@ -41,9 +42,9 @@ static ge::graphStatus InferShape(gert::InferShapeContext* context) } namespace ops { -class ArithprogressionCustom : public OpDef { +class ArangeCustom : public OpDef { public: - explicit ArithprogressionCustom(const char* name) : OpDef(name) + explicit ArangeCustom(const char* name) : OpDef(name) { this->Input("firstGm") .ParamType(REQUIRED) @@ -69,5 +70,5 @@ public: this->AICore().AddConfig("ascend910b"); } }; -OP_ADD(ArithprogressionCustom); +OP_ADD(ArangeCustom); } \ No newline at end of file diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom_tiling.h b/examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom_tiling.h similarity index 56% rename from examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom_tiling.h rename to examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom_tiling.h index 1daa71c35dece4521b50b942b150077747639e6f..251a3d970bb67a1ecc0d941d8f24a91863eed26a 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/op_host/arithprogression_custom_tiling.h +++ b/examples/index/arange/kernel_launch_method_by_framework/op_host/arange_custom_tiling.h @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -14,14 +15,14 @@ #include "tiling/tiling_api.h" namespace optiling { -BEGIN_TILING_DATA_DEF(ArithprogressionCustomTilingData) +BEGIN_TILING_DATA_DEF(ArangeCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, outSize); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(ArithprogressionCustom, ArithprogressionCustomTilingData) +REGISTER_TILING_DATA_CLASS(ArangeCustom, ArangeCustomTilingData) } // namespace optiling -void ComputeTiling(const uint32_t outSize, optiling::ArithprogressionCustomTilingData &tiling){ +void ComputeTiling(const uint32_t outSize, optiling::ArangeCustomTilingData &tiling){ tiling.set_outSize(outSize); } diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt similarity index 80% rename from examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt index c50a409a20bd0e0cce495824295a18799e4f8be1..6f85a73857c235c8439ddfaf819164d6b9a6b952 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt +++ b/examples/index/arange/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt @@ -1,4 +1,11 @@ -# set custom compile options +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx") add_ops_compile_options(ALL OPTIONS -g -O0) endif() diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/arithprogression_custom.cpp b/examples/index/arange/kernel_launch_method_by_framework/op_kernel/arange_custom.cpp similarity index 57% rename from examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/arithprogression_custom.cpp rename to examples/index/arange/kernel_launch_method_by_framework/op_kernel/arange_custom.cpp index 1ac97ad620858e00e916911900009ae85f49554b..ed8ffc1bb7e9942cf87517021a4f8dad57fcafb3 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/op_kernel/arithprogression_custom.cpp +++ b/examples/index/arange/kernel_launch_method_by_framework/op_kernel/arange_custom.cpp @@ -1,17 +1,18 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "kernel_operator.h" -#include "../../../../../../kernel_impl/arithprogression_kernel.h" +#include "../../../../../../kernel_impl/arange_kernel.h" -extern "C" __global__ __aicore__ void arithprogression_custom(GM_ADDR firstGm, GM_ADDR diffGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { +extern "C" __global__ __aicore__ void arange_custom(GM_ADDR firstGm, GM_ADDR diffGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { if ASCEND_IS_AIC { return; } @@ -20,7 +21,7 @@ extern "C" __global__ __aicore__ void arithprogression_custom(GM_ADDR firstGm, G auto first = reinterpret_cast<__gm__ float*>(firstGm); auto diff = reinterpret_cast<__gm__ float*>(diffGm); if (TILING_KEY_IS(1)) { - MyCustomKernel::KernelArithProgression op; + MyCustomKernel::KernelArange op; op.Init(dstGm, *first, *diff, vecTiling.outSize); op.Process(); } diff --git a/examples/index/arithprogression/kernel_launch_method_by_direct/scripts/gen_data.py b/examples/index/arange/kernel_launch_method_by_framework/scripts/gen_data.py similarity index 70% rename from examples/index/arithprogression/kernel_launch_method_by_direct/scripts/gen_data.py rename to examples/index/arange/kernel_launch_method_by_framework/scripts/gen_data.py index c10122ce59c08d6ee43132518d0e03cfb401af3d..11c542cc45ca8ae51e1d92ac62d0d0da07eeb138 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_direct/scripts/gen_data.py +++ b/examples/index/arange/kernel_launch_method_by_framework/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 -# Copyright (c) 2024 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import numpy as np import os diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/help.info b/examples/index/arange/kernel_launch_method_by_framework/scripts/help.info similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/scripts/help.info rename to examples/index/arange/kernel_launch_method_by_framework/scripts/help.info diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/install.sh b/examples/index/arange/kernel_launch_method_by_framework/scripts/install.sh similarity index 91% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/install.sh rename to examples/index/arange/kernel_launch_method_by_framework/scripts/install.sh index 8468c5a256f2c77fad5bf78ab108ca5b62aad672..48c93e67002c90b8b703c47596b6b49f2a2a006c 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/install.sh +++ b/examples/index/arange/kernel_launch_method_by_framework/scripts/install.sh @@ -258,7 +258,15 @@ if [ $? -ne 0 ];then exit 1 fi -# set the set_env.bash +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== if [ -n "${INSTALL_PATH}" ] && [ -d ${INSTALL_PATH} ]; then _ASCEND_CUSTOM_OPP_PATH=${targetdir}/${vendordir} bin_path="${_ASCEND_CUSTOM_OPP_PATH}/bin" diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/upgrade.sh b/examples/index/arange/kernel_launch_method_by_framework/scripts/upgrade.sh similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/scripts/upgrade.sh rename to examples/index/arange/kernel_launch_method_by_framework/scripts/upgrade.sh diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_framework/testcases/CMakeLists.txt similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/testcases/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_framework/testcases/CMakeLists.txt diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/cmake/fun.cmake b/examples/index/arange/kernel_launch_method_by_framework/testcases/cmake/fun.cmake similarity index 100% rename from examples/index/arithprogression/kernel_launch_method_by_framework/testcases/cmake/fun.cmake rename to examples/index/arange/kernel_launch_method_by_framework/testcases/cmake/fun.cmake diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt b/examples/index/arange/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt similarity index 67% rename from examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt rename to examples/index/arange/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt index 4841f092e6ec22c204b07e027a19f9ec3ee54145..3f6b3bd27ce91093a289b739abcd779e663ec34b 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt +++ b/examples/index/arange/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt @@ -1,10 +1,10 @@ -add_npu_target(OP ArithprogressionCustom SRC arithprogression_custom_main.cpp) +add_npu_target(OP ArangeCustom SRC arange_custom_main.cpp) add_custom_target(run_npu_test COMMAND echo "===============================================================================" COMMAND echo " Run NPU test at ${CMAKE_CURRENT_BINARY_DIR}" COMMAND echo "===============================================================================" - COMMAND $ + COMMAND $ COMMAND echo "===============================================================================" ) -add_dependencies(run_npu_test arithprogression_custom_npu) \ No newline at end of file +add_dependencies(run_npu_test arange_custom_npu) \ No newline at end of file diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/arithprogression_custom_main.cpp b/examples/index/arange/kernel_launch_method_by_framework/testcases/npu/arange_custom_main.cpp similarity index 87% rename from examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/arithprogression_custom_main.cpp rename to examples/index/arange/kernel_launch_method_by_framework/testcases/npu/arange_custom_main.cpp index 3623f937d95c32d271c5115338e7ffc94e22830b..e4173e0764271dd137a1e50109eee6cf5d60d7de 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/testcases/npu/arithprogression_custom_main.cpp +++ b/examples/index/arange/kernel_launch_method_by_framework/testcases/npu/arange_custom_main.cpp @@ -1,17 +1,18 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include #include #include "acl/acl_rt.h" #include "acl/acl.h" -#include "aclnn_arithprogression_custom.h" +#include "aclnn_arange_custom.h" #include "../../../../../common/data_utils.h" aclrtStream CreateStream(int32_t device) @@ -58,13 +59,13 @@ void DestroyTensor(aclTensor *tensors[], void *devMem[], int32_t tensorCount) } struct tensorInfo { - int64_t *dims; + int64_t* dims; int64_t dimCnt; aclDataType dtype; aclFormat fmt; }; -int64_t GetDataSize(struct tensorInfo *desc) +int64_t GetDataSize(struct tensorInfo* desc) { if (!desc->dims) { return 0; @@ -168,20 +169,20 @@ int32_t main(void) size_t workspaceSize = 0; aclOpExecutor *handle; - int32_t ret = aclnnArithprogressionCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[2], &workspaceSize, &handle); + int32_t ret = aclnnArangeCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[2], &workspaceSize, &handle); if (ret != ACL_SUCCESS) { - printf("aclnnArithprogressionCustomGetWorkspaceSize failed. error code is %u\n", ret); + printf("aclnnArangeCustomGetWorkspaceSize failed. error code is %u\n", ret); DestroyTensor(tensors, devMem, tensorCount); DestroyStream(stream, 0); return ret; } - printf("aclnnArithprogressionCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); + printf("aclnnArangeCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); void *workspace = nullptr; if (workspaceSize != 0) { CHECK_ACL(aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST)); } - ret = aclnnArithprogressionCustom(workspace, workspaceSize, handle, stream); - printf("aclnnArithprogressionCustom ret %u\n", ret); + ret = aclnnArangeCustom(workspace, workspaceSize, handle, stream); + printf("aclnnArangeCustom ret %u\n", ret); if (aclrtSynchronizeStreamWithTimeout(stream, TIMEOUT_MS) != ACL_SUCCESS) { printf("Synchronize stream failed\n"); } diff --git a/examples/readme.md b/examples/readme.md index e51ff5ea063a9fa67cc5d3ca70591c23d9eaf6cd..97e56fa52c6e9e5419b7bd6f861b126be44f28ec 100644 --- a/examples/readme.md +++ b/examples/readme.md @@ -213,7 +213,7 @@ - + @@ -243,17 +243,17 @@ - + - + - +
算子类型(OpType)ArithprogressionCustom
算子类型(OpType)ArangeCustom
算子输入
nameshapedata typeformat
transpose confusion_transpose transpose 对输入数据进行数据排布及Reshape操作。
utils init_global_memory fill 将Global Memory上的数据初始化为指定值。
index arithprogression arange 基于给定的起始值，等差值和长度，返回一个等差数列。
select selectwithbytesmask select 对输入srcTensor和srcScalar，根据maskTensor相应位置的值从二者中选取元素得到dstTensor。
diff --git a/examples/select/selectwithbytesmask/README.md b/examples/select/README.md similarity index 66% rename from examples/select/selectwithbytesmask/README.md rename to examples/select/README.md index fd6958945bd4f1fcb6affca0e2e6b5b464df47a8..4aced00becb0647ece3abb2ea4fbda8d832ba7de 100644 --- a/examples/select/selectwithbytesmask/README.md +++ b/examples/select/README.md @@ -2,13 +2,13 @@ ## 概述 -本样例介绍了调用SelectWithBytesMask高阶API实现SelectWithBytesMaskCustom单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 +本样例介绍了调用Select高阶API实现SelectCustom单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 -- 直调：使用核函数直调SelectWithBytesMaskCustom自定义算子。 +- 直调：使用核函数直调SelectCustom自定义算子。 核函数的基础调用方式，开发者完成算子核函数的开发和Tiling实现后，即可通过AscendCL运行时接口，完成算子的调用。 -- 框架调用：使用框架调用SelectWithBytesMaskCustom自定义算子。 +- 框架调用：使用框架调用SelectCustom自定义算子。 按照工程创建->算子实现->编译部署->算子调用的流程完成算子开发。整个过程都依赖于算子工程：基于工程代码框架完成算子核函数的开发和Tiling实现，通过工程编译脚本完成算子的编译部署，继而实现单算子调用或第三方框架中的算子调用。 @@ -17,7 +17,7 @@ | 调用方式 | 目录 | **描述** | | --------- | ------------------------------------------------------------ | ---------------------------------------------------------- | | 直调 | [kernel_launch_method_by_direct](./kernel_launch_method_by_direct) | host侧的核函数调用程序，包含CPU侧、NPU侧、仿真侧三种运行验证方法。 | -| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用SelectWithBytesMaskCustom算子。 | +| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用SelectCustom算子。 | ## 样例支持的产品型号为： - Atlas A2训练系列产品/Atlas 800I A2推理产品 @@ -34,12 +34,12 @@ ## 算子描述 -SelectWithBytesMaskCustom单算子，对输入src0Tensor和src1Scalar，根据maskTensor相应位置的值(非bit位)从二者中选取元素得到dstTensor。 +SelectCustom单算子，对输入src0Tensor和src1Scalar，根据maskTensor相应位置的值(非bit位)从二者中选取元素得到dstTensor。 -SelectWithBytesMaskCustom算子规格： +SelectCustom算子规格： - + @@ -51,19 +51,19 @@ SelectWithBytesMaskCustom算子规格： - +
算子类型(OpType)SelectWithBytesMaskCustom
算子类型(OpType)SelectCustom
算子输入
nameshapedata typeformat
dstGm2*32floatND
核函数名select_with_bytes_mask_custom
核函数名select_custom
## 算子实现介绍 -本样例中实现的是固定shape，输入src[2, 32]，mask[2, 32]，输出dst[2, 32]的SelectWithBytesMaskCustom算子。 +本样例中实现的是固定shape，输入src[2, 32]，mask[2, 32]，输出dst[2, 32]的SelectCustom算子。 - kernel实现 - 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用SelectWithBytesMask高阶API接口完成计算，得到最终结果，再搬出到外部存储上。 + 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Select高阶API接口完成计算，得到最终结果，再搬出到外部存储上。 - SelectWithBytesMaskCustom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行SelectWithBytesMask计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 + SelectCustom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Select计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 - tiling实现 - SelectWithBytesMaskCustom算子的tiling实现流程如下：首先获取SelectWithBytesMask接口能完成计算所需最大/最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小，然后根据SrcTensor和maskTensor的shape信息，确定其余tiling参数，最后将scalar类型的源操作数也包含在tiling中传递到kernel侧。 \ No newline at end of file + SelectCustom算子的tiling实现流程如下：首先获取Select接口能完成计算所需最大/最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小，然后根据SrcTensor和maskTensor的shape信息，确定其余tiling参数，最后将scalar类型的源操作数也包含在tiling中传递到kernel侧。 \ No newline at end of file diff --git a/examples/select/selectwithbytesmask/host_tiling/select_with_bytes_mask_custom_tiling.h b/examples/select/host_tiling/select_custom_tiling.h similarity index 76% rename from examples/select/selectwithbytesmask/host_tiling/select_with_bytes_mask_custom_tiling.h rename to examples/select/host_tiling/select_custom_tiling.h index 4892a75d0a930164c2238ccf0f2c21597a64655b..846e869ab918944e07faa22d0fd512f0f4d8465d 100644 --- a/examples/select/selectwithbytesmask/host_tiling/select_with_bytes_mask_custom_tiling.h +++ b/examples/select/host_tiling/select_custom_tiling.h @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -22,7 +23,7 @@ BEGIN_TILING_DATA_DEF(SelectCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, sharedTmpBufferSize); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(SelectWithBytesMaskCustom, SelectCustomTilingData) +REGISTER_TILING_DATA_CLASS(SelectCustom, SelectCustomTilingData) } // namespace optiling void ComputeTiling(const uint32_t firstAxis, const uint32_t srcSize, @@ -36,7 +37,7 @@ void ComputeTiling(const uint32_t firstAxis, const uint32_t srcSize, uint32_t maxValue = 0; uint32_t minValue = 0; uint32_t srcTypeSize = sizeof(float); - AscendC::GetSelectWithBytesMaskMaxMinTmpSize(src0Shape, src1Shape, srcTypeSize, maskShape, 1, false, maxValue, minValue); + AscendC::GetSelectMaxMinTmpSize(src0Shape, src1Shape, srcTypeSize, maskShape, 1, false, maxValue, minValue); uint32_t sharedTmpBufferSize = minValue; tiling.set_scalarValue(1.234f); tiling.set_firstAxis(firstAxis); diff --git a/examples/select/selectwithbytesmask/kernel_impl/select_with_bytes_mask_kernel.h b/examples/select/kernel_impl/select_kernel.h similarity index 85% rename from examples/select/selectwithbytesmask/kernel_impl/select_with_bytes_mask_kernel.h rename to examples/select/kernel_impl/select_kernel.h index b41749814d22b579770f15f2d9b6112ad32afa94..8c31c451d796e0180187071f49a31b212dae4b72 100644 --- a/examples/select/selectwithbytesmask/kernel_impl/select_with_bytes_mask_kernel.h +++ b/examples/select/kernel_impl/select_kernel.h @@ -1,15 +1,16 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_SELECT_WITH_BYTES_MASK_H -#define EXAMPLES_SELECT_WITH_BYTES_MASK_H +#ifndef EXAMPLES_SELECT_H +#define EXAMPLES_SELECT_H #include "kernel_operator.h" namespace MyCustomKernel { @@ -69,9 +70,9 @@ private: AscendC::LocalTensor tmpBuffer = sharedTmpBuffer.Get(); dstLocal = outQueue.AllocTensor(); - AscendC::SelectWithBytesMask(dstLocal, srcLocal, scalar, maskLocal, tmpBuffer, info); + AscendC::Select(dstLocal, srcLocal, scalar, maskLocal, tmpBuffer, info); // Reverse Select. - // AscendC::SelectWithBytesMask(dstLocal, scalar, srcLocal, maskLocal, tmpBuffer, info); + // AscendC::Select(dstLocal, scalar, srcLocal, maskLocal, tmpBuffer, info); outQueue.EnQue(dstLocal); maskQueue.FreeTensor(maskLocal); inQueueX.FreeTensor(srcLocal); @@ -101,4 +102,4 @@ private: srcType scalar = 0.0f; }; } -#endif // EXAMPLES_SELECT_WITH_BYTES_MASK_H \ No newline at end of file +#endif // EXAMPLES_SELECT_H \ No newline at end of file diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/CMakeLists.txt b/examples/select/kernel_launch_method_by_direct/CMakeLists.txt similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/CMakeLists.txt rename to examples/select/kernel_launch_method_by_direct/CMakeLists.txt diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/README.md b/examples/select/kernel_launch_method_by_direct/README.md similarity index 91% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/README.md rename to examples/select/kernel_launch_method_by_direct/README.md index 2768a7b3b254e572503a9f30b08c216cab26afe8..31056f50efc6a1306680e2178fd899018b88cfca 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/README.md +++ b/examples/select/kernel_launch_method_by_direct/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于Kernel直调算子工程，介绍了调用SelectWithBytesMask高阶API实现SelectWithBytesMaskCustom单算子，主要演示SelectWithBytesMask高阶API在Kernel直调工程中的调用。 +本样例基于Kernel直调算子工程，介绍了调用Select高阶API实现SelectCustom单算子，主要演示Select高阶API在Kernel直调工程中的调用。 ## 目录结构介绍 | 目录及文件 | 描述 | diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/cmake/cpu_lib.cmake b/examples/select/kernel_launch_method_by_direct/cmake/cpu_lib.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/cmake/cpu_lib.cmake rename to examples/select/kernel_launch_method_by_direct/cmake/cpu_lib.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/cmake/npu_lib.cmake b/examples/select/kernel_launch_method_by_direct/cmake/npu_lib.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/cmake/npu_lib.cmake rename to examples/select/kernel_launch_method_by_direct/cmake/npu_lib.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/main.cpp b/examples/select/kernel_launch_method_by_direct/main.cpp similarity index 89% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/main.cpp rename to examples/select/kernel_launch_method_by_direct/main.cpp index 261d17c2bc566677dcb941ab9710515fe0a89c37..02b59274518313fc0f8653b4601bad6fc335551e 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/main.cpp +++ b/examples/select/kernel_launch_method_by_direct/main.cpp @@ -1,21 +1,22 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "../../../common/data_utils.h" +#include "../../common/data_utils.h" #ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" -extern void select_with_bytes_mask_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *srcGm, +extern void select_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *srcGm, uint8_t *maskGm, uint8_t *dstGm, uint8_t *workspace, uint8_t *tiling); #else #include "tikicpulib.h" -extern "C" __global__ __aicore__ void select_with_bytes_mask_custom( +extern "C" __global__ __aicore__ void select_custom( GM_ADDR srcGm, GM_ADDR maskGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling); #endif @@ -91,7 +92,7 @@ int32_t main(int32_t argc, char *argv[]) memcpy_s(tiling, tilingFileSize, GenerateTiling(FIRSTAXIS, SRCLASTAXIS, MASKLASTAXIS), tilingFileSize); AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(select_with_bytes_mask_custom, blockDim, srcInput, maskInput, output, workspace, tiling); // use this macro for cpu debug + ICPU_RUN_KF(select_custom, blockDim, srcInput, maskInput, output, workspace, tiling); // use this macro for cpu debug WriteFile("../output/output.bin", output, srcSize); bool goldenResult = CompareResult(output, srcSize); @@ -138,7 +139,7 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtMemcpy(srcInputDevice, srcSize, srcInputHost, srcSize, ACL_MEMCPY_HOST_TO_DEVICE)); CHECK_ACL(aclrtMemcpy(maskInputDevice, maskSize, maskInputHost, maskSize, ACL_MEMCPY_HOST_TO_DEVICE)); - select_with_bytes_mask_custom_do(blockDim, nullptr, stream, srcInputDevice, maskInputDevice, outputDevice, workspaceDevice, tilingDevice); + select_custom_do(blockDim, nullptr, stream, srcInputDevice, maskInputDevice, outputDevice, workspaceDevice, tilingDevice); CHECK_ACL(aclrtSynchronizeStream(stream)); CHECK_ACL(aclrtMemcpy(outputHost, srcSize, outputDevice, srcSize, ACL_MEMCPY_DEVICE_TO_HOST)); diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/run.sh b/examples/select/kernel_launch_method_by_direct/run.sh similarity index 50% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/run.sh rename to examples/select/kernel_launch_method_by_direct/run.sh index a375bdd197b1ce5d3b71f2bec33604cdae5a717d..8979690e6b0c379d4e4d01ae10c61a863783d0d6 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/run.sh +++ b/examples/select/kernel_launch_method_by_direct/run.sh @@ -26,6 +26,21 @@ rm -rf build mkdir build cd build +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== +if [ "${RUN_MODE}" = "sim" ]; then + export LD_LIBRARY_PATH=$(echo $LD_LIBRARY_PATH | sed 's/\/.*\/runtime\/lib64://g') + export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/runtime/lib64/stub:$LD_LIBRARY_PATH +fi + +source $ASCEND_HOME_DIR/bin/setenv.bash export LD_LIBRARY_PATH=${ASCEND_HOME_DIR}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH set -euo pipefail diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/scripts/gen_data.py b/examples/select/kernel_launch_method_by_direct/scripts/gen_data.py similarity index 70% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/scripts/gen_data.py rename to examples/select/kernel_launch_method_by_direct/scripts/gen_data.py index a7ccf05826b831e7d3a3c754c3f3b1a54db79a72..0b96a50cefdf34f9ed24791ee8f2ace48f1bdf51 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/scripts/gen_data.py +++ b/examples/select/kernel_launch_method_by_direct/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 -# Copyright (c) 2024 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import numpy as np import os diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom.cpp b/examples/select/kernel_launch_method_by_direct/select_custom.cpp similarity index 68% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom.cpp rename to examples/select/kernel_launch_method_by_direct/select_custom.cpp index 26573fdb2b586206478f858593fed517f3c467ed..5add81f1bb626b1c8146446cbdf2c6f42fb6fae7 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom.cpp +++ b/examples/select/kernel_launch_method_by_direct/select_custom.cpp @@ -1,14 +1,15 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "kernel_operator.h" -#include "../kernel_impl/select_with_bytes_mask_kernel.h" +#include "../kernel_impl/select_kernel.h" __aicore__ inline void CopyTiling(MyCustomKernel::VecTiling* tiling, GM_ADDR tilingGM) { @@ -21,7 +22,7 @@ __aicore__ inline void CopyTiling(MyCustomKernel::VecTiling* tiling, GM_ADDR til return; } -extern "C" __global__ __aicore__ void select_with_bytes_mask_custom( +extern "C" __global__ __aicore__ void select_custom( GM_ADDR srcGm, GM_ADDR maskGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { MyCustomKernel::VecTiling tilingData; @@ -41,9 +42,9 @@ extern "C" __global__ __aicore__ void select_with_bytes_mask_custom( #ifndef ASCENDC_CPU_DEBUG // call of kernel function -void select_with_bytes_mask_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *srcGm, +void select_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *srcGm, uint8_t *maskGm, uint8_t *dstGm, uint8_t *workspace, uint8_t *tiling) { - select_with_bytes_mask_custom<<>>(srcGm, maskGm, dstGm, workspace, tiling); + select_custom<<>>(srcGm, maskGm, dstGm, workspace, tiling); } #endif \ No newline at end of file diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom_tiling.cpp b/examples/select/kernel_launch_method_by_direct/select_custom_tiling.cpp similarity index 70% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom_tiling.cpp rename to examples/select/kernel_launch_method_by_direct/select_custom_tiling.cpp index f44349b287f58501c1cfde8411528e6014eb2762..94541cf474529e90c08e2e66ab5632abfb5ceec1 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_direct/select_custom_tiling.cpp +++ b/examples/select/kernel_launch_method_by_direct/select_custom_tiling.cpp @@ -1,10 +1,11 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -14,7 +15,7 @@ #include #include #include "tiling/tiling_api.h" -#include "../host_tiling/select_with_bytes_mask_custom_tiling.h" +#include "../host_tiling/select_custom_tiling.h" uint8_t *GetTilingBuf(optiling::SelectCustomTilingData *tilingData) { uint32_t tilingSize = sizeof(optiling::SelectCustomTilingData); diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/CMakeLists.txt b/examples/select/kernel_launch_method_by_framework/CMakeLists.txt similarity index 81% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/CMakeLists.txt rename to examples/select/kernel_launch_method_by_framework/CMakeLists.txt index 584132d80993d309434fb1303de83910a1989aba..83a789d3366aa7794f09d01ba16d73c0aeac7f69 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/CMakeLists.txt +++ b/examples/select/kernel_launch_method_by_framework/CMakeLists.txt @@ -33,7 +33,14 @@ if(ENABLE_TEST AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/testcases) add_subdirectory(testcases) endif() -# modify vendor_name in install.sh and upgrade.sh +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/scripts/install.sh ${CMAKE_BINARY_DIR}/scripts/upgrade.sh COMMAND mkdir -p ${CMAKE_BINARY_DIR}/scripts COMMAND cp -r ${CMAKE_SOURCE_DIR}/scripts/* ${CMAKE_BINARY_DIR}/scripts/ diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/CMakePresets.json b/examples/select/kernel_launch_method_by_framework/CMakePresets.json similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/CMakePresets.json rename to examples/select/kernel_launch_method_by_framework/CMakePresets.json diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/README.md b/examples/select/kernel_launch_method_by_framework/README.md similarity index 95% rename from examples/index/arithprogression/kernel_launch_method_by_framework/README.md rename to examples/select/kernel_launch_method_by_framework/README.md index 829d5a3009cd7e5164401a0f27de781ea72d80b0..b244a23eb72ded0fe1465c2724a10abc899a1dcb 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/README.md +++ b/examples/select/kernel_launch_method_by_framework/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于自定义算子工程，介绍了调用ArithProgression高阶API实现arithprogression单算子，主要演示ArithProgression高阶API在自定义算子工程中的调用。 +本样例基于自定义算子工程，介绍了调用Select高阶API实现SelectCustom单算子，主要演示Select高阶API在自定义算子工程中的调用。 ## 目录结构 | 目录 | 描述 | @@ -75,7 +75,7 @@ cd build_out ### 5.执行样例 在build_out目录下执行如下命令 ``` -./arithprogression_custom_npu +./select_custom_npu ``` ### 6.NPU仿真模式运行（可选） 若要执行NPU仿真，在build_out目录下执行如下命令： diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/build.sh b/examples/select/kernel_launch_method_by_framework/build.sh similarity index 95% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/build.sh rename to examples/select/kernel_launch_method_by_framework/build.sh index 0813af07bbc6da272e38b7019e81b538bbbf8b30..229910e0133e72beadd6055c3ebb269442a6ea40 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/build.sh +++ b/examples/select/kernel_launch_method_by_framework/build.sh @@ -3,7 +3,6 @@ script_path=$(realpath $(dirname $0)) source $ASCEND_HOME_DIR/bin/setenv.bash cp -rf ../host_tiling/* op_host/ -rm -rf ./cmake/util ln -s $ASCEND_HOME_DIR/tools/op_project_templates/ascendc/customize/cmake/util/ ./cmake/util mkdir -p build_out rm -rf build_out/* @@ -70,8 +69,3 @@ else fi -# for debug -# cd build_out -# make -# cpack -# verbose append -v \ No newline at end of file diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/config.cmake b/examples/select/kernel_launch_method_by_framework/cmake/config.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/config.cmake rename to examples/select/kernel_launch_method_by_framework/cmake/config.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/func.cmake b/examples/select/kernel_launch_method_by_framework/cmake/func.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/func.cmake rename to examples/select/kernel_launch_method_by_framework/cmake/func.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/intf.cmake b/examples/select/kernel_launch_method_by_framework/cmake/intf.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/intf.cmake rename to examples/select/kernel_launch_method_by_framework/cmake/intf.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/makeself.cmake b/examples/select/kernel_launch_method_by_framework/cmake/makeself.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/makeself.cmake rename to examples/select/kernel_launch_method_by_framework/cmake/makeself.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/util b/examples/select/kernel_launch_method_by_framework/cmake/util similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/cmake/util rename to examples/select/kernel_launch_method_by_framework/cmake/util diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_host/CMakeLists.txt b/examples/select/kernel_launch_method_by_framework/op_host/CMakeLists.txt similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_host/CMakeLists.txt rename to examples/select/kernel_launch_method_by_framework/op_host/CMakeLists.txt diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_host/select_with_bytes_mask_custom.cpp b/examples/select/kernel_launch_method_by_framework/op_host/select_custom.cpp similarity index 82% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_host/select_with_bytes_mask_custom.cpp rename to examples/select/kernel_launch_method_by_framework/op_host/select_custom.cpp index 0c9cb2cf88b6c909070149523fad8370f5f7a8cf..45a151a75272b4878f7d08bb2a486d2df1fdf011 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_host/select_with_bytes_mask_custom.cpp +++ b/examples/select/kernel_launch_method_by_framework/op_host/select_custom.cpp @@ -1,14 +1,15 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "select_with_bytes_mask_custom_tiling.h" +#include "select_custom_tiling.h" #include "register/op_def_registry.h" namespace optiling { @@ -43,9 +44,9 @@ static ge::graphStatus InferShape(gert::InferShapeContext* context) namespace ops { -class SelectWithBytesMaskCustom : public OpDef { +class SelectCustom : public OpDef { public: - explicit SelectWithBytesMaskCustom(const char* name) : OpDef(name) + explicit SelectCustom(const char* name) : OpDef(name) { this->Input("srcGm") .ParamType(REQUIRED) @@ -71,5 +72,5 @@ public: this->AICore().AddConfig("ascend910b"); } }; -OP_ADD(SelectWithBytesMaskCustom); +OP_ADD(SelectCustom); } diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt b/examples/select/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt similarity index 80% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt rename to examples/select/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt index c50a409a20bd0e0cce495824295a18799e4f8be1..6f85a73857c235c8439ddfaf819164d6b9a6b952 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt +++ b/examples/select/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt @@ -1,4 +1,11 @@ -# set custom compile options +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx") add_ops_compile_options(ALL OPTIONS -g -O0) endif() diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/select_with_bytes_mask_custom.cpp b/examples/select/kernel_launch_method_by_framework/op_kernel/select_custom.cpp similarity index 61% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/select_with_bytes_mask_custom.cpp rename to examples/select/kernel_launch_method_by_framework/op_kernel/select_custom.cpp index f25d09ca1d3ebbcc01323c468ca50f0411197193..eec6d1a68fa3287530d86cccf39073eaa805ce7a 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/op_kernel/select_with_bytes_mask_custom.cpp +++ b/examples/select/kernel_launch_method_by_framework/op_kernel/select_custom.cpp @@ -1,17 +1,18 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "kernel_operator.h" -#include "../../../../../../kernel_impl/select_with_bytes_mask_kernel.h" +#include "../../../../../../kernel_impl/select_kernel.h" -extern "C" __global__ __aicore__ void select_with_bytes_mask_custom(GM_ADDR srcGm, GM_ADDR maskGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { +extern "C" __global__ __aicore__ void select_custom(GM_ADDR srcGm, GM_ADDR maskGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { if ASCEND_IS_AIC { return; } diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/gen_data.py b/examples/select/kernel_launch_method_by_framework/scripts/gen_data.py similarity index 70% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/gen_data.py rename to examples/select/kernel_launch_method_by_framework/scripts/gen_data.py index a7ccf05826b831e7d3a3c754c3f3b1a54db79a72..0b96a50cefdf34f9ed24791ee8f2ace48f1bdf51 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/gen_data.py +++ b/examples/select/kernel_launch_method_by_framework/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 -# Copyright (c) 2024 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import numpy as np import os diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/help.info b/examples/select/kernel_launch_method_by_framework/scripts/help.info similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/help.info rename to examples/select/kernel_launch_method_by_framework/scripts/help.info diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/install.sh b/examples/select/kernel_launch_method_by_framework/scripts/install.sh similarity index 91% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/install.sh rename to examples/select/kernel_launch_method_by_framework/scripts/install.sh index 8468c5a256f2c77fad5bf78ab108ca5b62aad672..48c93e67002c90b8b703c47596b6b49f2a2a006c 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/install.sh +++ b/examples/select/kernel_launch_method_by_framework/scripts/install.sh @@ -258,7 +258,15 @@ if [ $? -ne 0 ];then exit 1 fi -# set the set_env.bash +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== if [ -n "${INSTALL_PATH}" ] && [ -d ${INSTALL_PATH} ]; then _ASCEND_CUSTOM_OPP_PATH=${targetdir}/${vendordir} bin_path="${_ASCEND_CUSTOM_OPP_PATH}/bin" diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/upgrade.sh b/examples/select/kernel_launch_method_by_framework/scripts/upgrade.sh similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/scripts/upgrade.sh rename to examples/select/kernel_launch_method_by_framework/scripts/upgrade.sh diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/CMakeLists.txt b/examples/select/kernel_launch_method_by_framework/testcases/CMakeLists.txt similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/CMakeLists.txt rename to examples/select/kernel_launch_method_by_framework/testcases/CMakeLists.txt diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/cmake/fun.cmake b/examples/select/kernel_launch_method_by_framework/testcases/cmake/fun.cmake similarity index 100% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/cmake/fun.cmake rename to examples/select/kernel_launch_method_by_framework/testcases/cmake/fun.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt b/examples/select/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt similarity index 66% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt rename to examples/select/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt index ab80a7da55b5655e2e6214314c43fe47e82b815d..d3188a793834a7bfa0431419df31022273ece5fe 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt +++ b/examples/select/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt @@ -1,10 +1,10 @@ -add_npu_target(OP InitGlobalMemoryCustom SRC init_global_memory_custom_main.cpp) +add_npu_target(OP SelectCustom SRC select_custom_main.cpp) add_custom_target(run_npu_test COMMAND echo "===============================================================================" COMMAND echo " Run NPU test at ${CMAKE_CURRENT_BINARY_DIR}" COMMAND echo "===============================================================================" - COMMAND $ + COMMAND $ COMMAND echo "===============================================================================" ) -add_dependencies(run_npu_test init_global_memory_custom_npu) \ No newline at end of file +add_dependencies(run_npu_test select_custom_npu) \ No newline at end of file diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/select_with_bytes_mask_custom_main.cpp b/examples/select/kernel_launch_method_by_framework/testcases/npu/select_custom_main.cpp similarity index 87% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/select_with_bytes_mask_custom_main.cpp rename to examples/select/kernel_launch_method_by_framework/testcases/npu/select_custom_main.cpp index 0a2ea707025b2d7389edf496b5510e49a7b548af..52b0b79252270d2b43b5be107d858186404b6dab 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/select_with_bytes_mask_custom_main.cpp +++ b/examples/select/kernel_launch_method_by_framework/testcases/npu/select_custom_main.cpp @@ -1,18 +1,19 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. +/* + * This program is free software, you can redistribute it and/or modify it. + * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include #include #include "acl/acl_rt.h" #include "acl/acl.h" -#include "aclnn_select_with_bytes_mask_custom.h" -#include "../../../../../common/data_utils.h" +#include "aclnn_select_custom.h" +#include "../../../../common/data_utils.h" aclrtStream CreateStream(int32_t device) { @@ -58,7 +59,7 @@ void DestroyTensor(aclTensor *tensors[], void *devMem[], int32_t tensorCount) } struct tensorInfo { - int64_t *dims; + int64_t* dims; int64_t dimCnt; aclDataType dtype; aclFormat fmt; @@ -161,20 +162,20 @@ int32_t main(void) size_t workspaceSize = 0; aclOpExecutor *handle; - int32_t ret = aclnnSelectWithBytesMaskCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[2], &workspaceSize, &handle); + int32_t ret = aclnnSelectCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[2], &workspaceSize, &handle); if (ret != ACL_SUCCESS) { - printf("aclnnSelectWithBytesMaskCustomGetWorkspaceSize failed. error code is %u\n", ret); + printf("aclnnSelectCustomGetWorkspaceSize failed. error code is %u\n", ret); DestroyTensor(tensors, devMem, tensorCount); DestroyStream(stream, 0); return ret; } - printf("aclnnSelectWithBytesMaskCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); + printf("aclnnSelectCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); void *workspace = nullptr; if (workspaceSize != 0) { CHECK_ACL(aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST)); } - ret = aclnnSelectWithBytesMaskCustom(workspace, workspaceSize, handle, stream); - printf("aclnnSelectWithBytesMaskCustom ret %u\n", ret); + ret = aclnnSelectCustom(workspace, workspaceSize, handle, stream); + printf("aclnnSelectCustom ret %u\n", ret); if (aclrtSynchronizeStreamWithTimeout(stream, TIMEOUT_MS) != ACL_SUCCESS) { printf("Synchronize stream failed\n"); } diff --git a/examples/transpose/confusion_transpose/CMakeLists.txt b/examples/transpose/CMakeLists.txt similarity index 73% rename from examples/transpose/confusion_transpose/CMakeLists.txt rename to examples/transpose/CMakeLists.txt index ef7fd4a9d807536ce598ce80ee962018ad37504e..7764f8d2f2b3ef645f17ce63753001dc42ee0d5b 100644 --- a/examples/transpose/confusion_transpose/CMakeLists.txt +++ b/examples/transpose/CMakeLists.txt @@ -1,11 +1,11 @@ +# This program is free software, you can redistribute it and/or modify it. # Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== cmake_minimum_required(VERSION 3.16) project(Ascend_c) @@ -27,7 +27,7 @@ if(CMAKE_INSTALL_PREFIX STREQUAL /usr/local) endif() file(GLOB KERNEL_FILES - ${CMAKE_CURRENT_SOURCE_DIR}/op_kernel/confusion_transpose_custom.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/op_kernel/transpose_custom.cpp ) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") @@ -39,29 +39,29 @@ else() message("invalid RUN_MODE: ${RUN_MODE}") endif() -add_executable(confusion_transpose_direct_kernel_op +add_executable(transpose_direct_kernel_op ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/op_host/confusion_transpose_custom_tiling.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/op_host/transpose_custom_tiling.cpp ) -target_compile_options(confusion_transpose_direct_kernel_op PRIVATE +target_compile_options(transpose_direct_kernel_op PRIVATE $:-g>> -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 ) -target_compile_definitions(confusion_transpose_direct_kernel_op PRIVATE +target_compile_definitions(transpose_direct_kernel_op PRIVATE $<$>:CUSTOM_ASCEND310P> ) -target_include_directories(confusion_transpose_direct_kernel_op PRIVATE +target_include_directories(transpose_direct_kernel_op PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} $:${ASCEND_CANN_PACKAGE_PATH}/include>> $:${ASCEND_CANN_PACKAGE_PATH}/runtime/include>> ) -target_link_libraries(confusion_transpose_direct_kernel_op PRIVATE +target_link_libraries(transpose_direct_kernel_op PRIVATE $,$>:host_intf_pub>> $:tikicpulib::${SOC_VERSION}>> $:ascendcl>> @@ -75,7 +75,7 @@ target_link_libraries(confusion_transpose_direct_kernel_op PRIVATE graph_base ) -install(TARGETS confusion_transpose_direct_kernel_op +install(TARGETS transpose_direct_kernel_op LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} diff --git a/examples/transpose/confusion_transpose/README.md b/examples/transpose/README.md similarity index 76% rename from examples/transpose/confusion_transpose/README.md rename to examples/transpose/README.md index 478c65bfc10f11ab43e6c97a6c414337d47740b2..cb6dd3731a190f79fc9b343b3d2c71a33090aa6d 100644 --- a/examples/transpose/confusion_transpose/README.md +++ b/examples/transpose/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例介绍了调用ConfusionTranspose高阶API实现confusionTranspose算子，并按照核函数直调的方式分别给出了对应的端到端实现，关于ConfusionTranspose高阶API的具体内容请参考《[Ascend C算子开发接口](https://hiascend.com/document/redirect/CannCommunityAscendCApi)》中的 "高阶 API > 变形 > ConfusionTranspose" 章节。 +本样例介绍了调用Transpose高阶API实现Transpose算子，并按照核函数直调的方式分别给出了对应的端到端实现，关于Transpose高阶API的具体内容请参考《[Ascend C算子开发接口](https://hiascend.com/document/redirect/CannCommunityAscendCApi)》中的 "高阶 API > 变形 > Transpose" 章节。 本样例以直调的方式调用算子核函数。 @@ -27,12 +27,12 @@ ## 算子描述 - 算子功能 - ConfusionTransposeCustom算子，对输入数据进行数据排布及Reshape操作。 + TransposeCustom算子，对输入数据进行数据排布及Reshape操作。 - 算子规格： - + @@ -40,12 +40,12 @@ - +
算子类型(OpType)ConfusionTransposeCustom
算子类型(OpType)TransposeCustom
算子输入
nameshapedata typeformat
算子输出
y1*64*2*32halfND
核函数名confusion_transpose_custom
核函数名transpose_custom
## 算子实现介绍 -本样例中实现的是固定shape为输入x[1, 2, 2, 4, 16, 16]，输出y[1, 64, 2, 32]的ConfusionTransposeCustom算子，针对NZ2ND场景，实现1、2轴互换。 +本样例中实现的是固定shape为输入x[1, 2, 2, 4, 16, 16]，输出y[1, 64, 2, 32]的TransposeCustom算子，针对NZ2ND场景，实现1、2轴互换。 输入Tensor { shape:[B, N, H/N/16, S/16, 16, 16], origin_shape:[B, N, S, H/N], format:"NZ", origin_format:"ND"} @@ -53,13 +53,13 @@ - Kernel实现 - 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用ConfusionTranspose高阶API接口完成confusionTranspose计算，得到最终结果，再搬出到外部存储上。 + 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Transpose高阶API接口完成Transpose计算，得到最终结果，再搬出到外部存储上。 - ConfusionTransposeCustom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行confusionTranspose计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 + TransposeCustom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Transpose计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 - Tiling实现 - 使用Ascend C提供的GetConfusionTransposeTilingInfo接口，获取所需的Tiling参数，并调用GetConfusionTransposeMaxMinTmpSize接口获取ConfusionTranspose接口计算所需的临时空间大小。 + 使用Ascend C提供的GetTransposeTilingInfo接口，获取所需的Tiling参数，并调用GetTransposeMaxMinTmpSize接口获取Transpose接口计算所需的临时空间大小。 ## 编译运行样例 diff --git a/examples/transpose/confusion_transpose/cmake/cpu_lib.cmake b/examples/transpose/cmake/cpu_lib.cmake similarity index 100% rename from examples/transpose/confusion_transpose/cmake/cpu_lib.cmake rename to examples/transpose/cmake/cpu_lib.cmake diff --git a/examples/transpose/confusion_transpose/cmake/npu_lib.cmake b/examples/transpose/cmake/npu_lib.cmake similarity index 100% rename from examples/transpose/confusion_transpose/cmake/npu_lib.cmake rename to examples/transpose/cmake/npu_lib.cmake diff --git a/examples/transpose/confusion_transpose/main.cpp b/examples/transpose/main.cpp similarity index 85% rename from examples/transpose/confusion_transpose/main.cpp rename to examples/transpose/main.cpp index 8bd7759de4eca059b07aa5944a13ff67b6ffbfee..aff795e4bed703dc3af972876dfd797db4662e44 100644 --- a/examples/transpose/confusion_transpose/main.cpp +++ b/examples/transpose/main.cpp @@ -1,21 +1,22 @@ /* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "../../common/data_utils.h" +#include "../common/data_utils.h" #ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" -#include "aclrtlaunch_confusion_transpose_custom.h" +#include "aclrtlaunch_transpose_custom.h" #include "tiling/platform/platform_ascendc.h" #else #include "tikicpulib.h" -extern "C" __global__ __aicore__ void confusion_transpose_custom(GM_ADDR x, GM_ADDR y, GM_ADDR tiling); +extern "C" __global__ __aicore__ void transpose_custom(GM_ADDR x, GM_ADDR y, GM_ADDR tiling); #endif namespace { @@ -46,7 +47,7 @@ int32_t main(int32_t argc, char *argv[]) { memcpy_s(tiling, tilingSize, buf, tilingSize); AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(confusion_transpose_custom, USED_CORE_NUM, x, y, tiling); + ICPU_RUN_KF(transpose_custom, USED_CORE_NUM, x, y, tiling); WriteFile("../output/output.bin", y, outputSize); @@ -79,7 +80,7 @@ int32_t main(int32_t argc, char *argv[]) { CHECK_ACL(aclrtMemcpy(tilingDevice, tilingSize, buf, tilingSize, ACL_MEMCPY_HOST_TO_DEVICE)); // Execute the kernel - ACLRT_LAUNCH_KERNEL(confusion_transpose_custom) + ACLRT_LAUNCH_KERNEL(transpose_custom) (USED_CORE_NUM, stream, xDevice, yDevice, tilingDevice); // Wait for the stop event to complete diff --git a/examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.cpp b/examples/transpose/op_host/transpose_custom_tiling.cpp similarity index 55% rename from examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.cpp rename to examples/transpose/op_host/transpose_custom_tiling.cpp index cfb1244e4123eb2f522c6142dd453cb270169ff8..c5d53ecdc5cf22259f56bd8d4b8da406fd9d3202 100644 --- a/examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.cpp +++ b/examples/transpose/op_host/transpose_custom_tiling.cpp @@ -1,24 +1,25 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "confusion_transpose_custom_tiling.h" +#include "transpose_custom_tiling.h" -uint8_t *GetTilingBuf(optiling::ConfusionTransposeCustomTilingData *tilingData) { - uint32_t tilingSize = sizeof(optiling::ConfusionTransposeCustomTilingData); +uint8_t *GetTilingBuf(optiling::TransposeCustomTilingData *tilingData) { + uint32_t tilingSize = sizeof(optiling::TransposeCustomTilingData); uint8_t *buf = (uint8_t *)malloc(tilingSize); tilingData->SaveToBuffer(buf, tilingSize); return buf; } uint8_t *GenerateTiling(uint32_t b, uint32_t n, uint32_t s, uint32_t hnDiv) { - optiling::ConfusionTransposeCustomTilingData tiling; + optiling::TransposeCustomTilingData tiling; ComputeTiling(b, n, s, hnDiv, tiling); return GetTilingBuf(&tiling); } \ No newline at end of file diff --git a/examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.h b/examples/transpose/op_host/transpose_custom_tiling.h similarity index 54% rename from examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.h rename to examples/transpose/op_host/transpose_custom_tiling.h index b5098bf30ba2e001945de42436093ff2e9c13b2f..933caccc8ab643ca8cf0b2b24354c72a5177deaa 100644 --- a/examples/transpose/confusion_transpose/op_host/confusion_transpose_custom_tiling.h +++ b/examples/transpose/op_host/transpose_custom_tiling.h @@ -1,32 +1,33 @@ /* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_TRANSPOSE_CONFUSION_TRANSPOSE_CUSTOM_TILING_H -#define EXAMPLES_TRANSPOSE_CONFUSION_TRANSPOSE_CUSTOM_TILING_H +#ifndef EXAMPLES_TRANSPOSE_CUSTOM_TILING_H +#define EXAMPLES_TRANSPOSE_CUSTOM_TILING_H #include "register/tilingdata_base.h" #include "tiling/platform/platform_ascendc.h" #include "tiling/tiling_api.h" namespace optiling { -BEGIN_TILING_DATA_DEF(ConfusionTransposeCustomTilingData) +BEGIN_TILING_DATA_DEF(TransposeCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, b); TILING_DATA_FIELD_DEF(uint32_t, n); TILING_DATA_FIELD_DEF(uint32_t, s); TILING_DATA_FIELD_DEF(uint32_t, hnDiv); - TILING_DATA_FIELD_DEF_STRUCT(ConfusionTransposeTiling, confusionTransposeTilingData); + TILING_DATA_FIELD_DEF_STRUCT(ConfusionTransposeTiling, ConfusionTransposeTilingData); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(ConfusionTransposeCustom, ConfusionTransposeCustomTilingData) +REGISTER_TILING_DATA_CLASS(TransposeCustom, TransposeCustomTilingData) } // namespace optiling void ComputeTiling(uint32_t b, uint32_t n, uint32_t s, uint32_t hnDiv, - optiling::ConfusionTransposeCustomTilingData &tiling) { + optiling::TransposeCustomTilingData &tiling) { tiling.set_b(b); tiling.set_n(n); tiling.set_s(s); @@ -37,9 +38,9 @@ void ComputeTiling(uint32_t b, uint32_t n, uint32_t s, uint32_t hnDiv, uint32_t maxValue = 0; uint32_t minValue = 0; - AscendC::GetConfusionTransposeMaxMinTmpSize(srcShape, sizeof(uint16_t), 1, maxValue, minValue); + AscendC::GetTransposeMaxMinTmpSize(srcShape, sizeof(uint16_t), 1, maxValue, minValue); const uint32_t stackBufferSize = minValue; - AscendC::GetConfusionTransposeTilingInfo(srcShape, stackBufferSize, sizeof(uint16_t), 1, tiling.confusionTransposeTilingData); + AscendC::GetTransposeTilingInfo(srcShape, stackBufferSize, sizeof(uint16_t), 1, tiling.ConfusionTransposeTilingData); } -#endif // EXAMPLES_TRANSPOSE_CONFUSION_TRANSPOSE_CUSTOM_TILING_H +#endif // EXAMPLES_TRANSPOSE_CUSTOM_TILING_H diff --git a/examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom.cpp b/examples/transpose/op_kernel/transpose_custom.cpp similarity index 67% rename from examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom.cpp rename to examples/transpose/op_kernel/transpose_custom.cpp index fd6ff38a3a54a5058b24a3d4a8a7671e638bd5e6..dd302a9a6489373c0561476607fdb4fc8de97830 100644 --- a/examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom.cpp +++ b/examples/transpose/op_kernel/transpose_custom.cpp @@ -1,13 +1,14 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "confusion_transpose_custom_impl.h" +#include "transpose_custom_impl.h" #include "kernel_operator.h" __aicore__ inline void CopyTiling(MyCustomKernel::VecTiling *tiling, GM_ADDR tilingGM) { @@ -20,12 +21,12 @@ __aicore__ inline void CopyTiling(MyCustomKernel::VecTiling *tiling, GM_ADDR til return; } -extern "C" __global__ __aicore__ void confusion_transpose_custom(GM_ADDR x, GM_ADDR y, GM_ADDR tiling) { +extern "C" __global__ __aicore__ void transpose_custom(GM_ADDR x, GM_ADDR y, GM_ADDR tiling) { if ASCEND_IS_AIC { return; } AscendC::TPipe pipe; - MyCustomKernel::KernelConfusionTranspose op; + MyCustomKernel::KernelTranspose op; MyCustomKernel::VecTiling tilingData; CopyTiling(&tilingData, tiling); op.Init(x, y, tilingData, &pipe); diff --git a/examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom_impl.h b/examples/transpose/op_kernel/transpose_custom_impl.h similarity index 79% rename from examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom_impl.h rename to examples/transpose/op_kernel/transpose_custom_impl.h index ae2ea6ac167c69077365b5676c06c4edc41cb55b..d66bc8b43d266f1a6c3e7956a30d361b1d527f16 100644 --- a/examples/transpose/confusion_transpose/op_kernel/confusion_transpose_custom_impl.h +++ b/examples/transpose/op_kernel/transpose_custom_impl.h @@ -1,15 +1,16 @@ /* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_TRANSPOSE_CONFUSION_TRANSPOSE_CUSTOM_IMPL_H -#define EXAMPLES_TRANSPOSE_CONFUSION_TRANSPOSE_CUSTOM_IMPL_H +#ifndef EXAMPLES_TRANSPOSE_CUSTOM_IMPL_H +#define EXAMPLES_TRANSPOSE_CUSTOM_IMPL_H #include "kernel_operator.h" @@ -19,13 +20,13 @@ struct VecTiling { uint32_t n; uint32_t s; uint32_t hnDiv; - ConfusionTransposeTiling confusionTransposeTilingData; + ConfusionTransposeTiling ConfusionTransposeTilingData; }; template -class KernelConfusionTranspose { +class KernelTranspose { public: - __aicore__ inline KernelConfusionTranspose() {} + __aicore__ inline KernelTranspose() {} __aicore__ inline void Init(__gm__ uint8_t *srcGm, __gm__ uint8_t *dstGm, const VecTiling &tilingData, AscendC::TPipe *pipeIn) { ASCENDC_ASSERT(AscendC::GetBlockNum() != 0, { KERNEL_LOG(KERNEL_ERROR, "block dim can not be zero!"); }); this->b = tilingData.b; @@ -39,7 +40,7 @@ public: pipe = pipeIn; pipe->InitBuffer(inQueueSrcVecIn, 1, b * n * s * hnDiv * sizeof(T)); pipe->InitBuffer(inQueueSrcVecOut, 1, b * n * s * hnDiv * sizeof(T)); - this->tiling = tilingData.confusionTransposeTilingData; + this->tiling = tilingData.ConfusionTransposeTilingData; } __aicore__ inline void Process() { CopyIn(); @@ -56,7 +57,7 @@ private: __aicore__ inline void Compute() { AscendC::LocalTensor srcLocal = inQueueSrcVecIn.DeQue(); AscendC::LocalTensor dstLocal = inQueueSrcVecOut.AllocTensor(); - AscendC::ConfusionTranspose(dstLocal, srcLocal, AscendC::TransposeType::TRANSPOSE_NZ2ND_0213, this->tiling); + AscendC::Transpose(dstLocal, srcLocal, AscendC::TransposeType::TRANSPOSE_NZ2ND_0213, this->tiling); inQueueSrcVecOut.EnQue(dstLocal); inQueueSrcVecIn.FreeTensor(srcLocal); } diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/run.sh b/examples/transpose/run.sh similarity index 41% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/run.sh rename to examples/transpose/run.sh index fa109954b5ed8bb5a2b837eaaa22cfaa2e82dfdb..28415f63be3febbc108608f5a4a15fa493e119ba 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/run.sh +++ b/examples/transpose/run.sh @@ -1,5 +1,15 @@ #!/bin/bash +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== + SHORT=r:,v:, LONG=run-mode:,soc-version:, OPTS=$(getopt -a --options $SHORT --longoptions $LONG -- "$@") @@ -26,18 +36,27 @@ rm -rf build mkdir build cd build +# in case of running op in simulator, use stub so instead +if [ "${RUN_MODE}" = "sim" ]; then + export LD_LIBRARY_PATH=$(echo $LD_LIBRARY_PATH | sed 's/\/.*\/runtime\/lib64://g') + export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/runtime/lib64/stub:$LD_LIBRARY_PATH +fi + +source $ASCEND_HOME_DIR/bin/setenv.bash export LD_LIBRARY_PATH=${ASCEND_HOME_DIR}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH -set -euo pipefail cmake -DRUN_MODE=${RUN_MODE} -DSOC_VERSION=${SOC_VERSION} -DASCEND_CANN_PACKAGE_PATH=${ASCEND_HOME_DIR} .. make -j16 if [ "${RUN_MODE}" = "npu" ]; then - ./init_global_memory_direct_kernel_op + ./transpose_direct_kernel_op elif [ "${RUN_MODE}" = "sim" ]; then export ASCEND_TOOLKIT_HOME=${ASCEND_HOME_DIR} export ASCEND_HOME_PATH=${ASCEND_HOME_DIR} - msprof op simulator --application=./init_global_memory_direct_kernel_op + msprof op simulator --application=./transpose_direct_kernel_op elif [ "${RUN_MODE}" = "cpu" ]; then - ./init_global_memory_direct_kernel_op + ./transpose_direct_kernel_op fi + +cd .. +python3 scripts/verify_data.py output/output.bin output/golden.bin diff --git a/examples/transpose/confusion_transpose/scripts/gen_data.py b/examples/transpose/scripts/gen_data.py similarity index 77% rename from examples/transpose/confusion_transpose/scripts/gen_data.py rename to examples/transpose/scripts/gen_data.py index 7afe9444f5ec4ee510fdcf4edccd41bd777b7771..f681fe6ce1cb4d44d0f819e53b0e9ad8401881d7 100644 --- a/examples/transpose/confusion_transpose/scripts/gen_data.py +++ b/examples/transpose/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 +# This program is free software, you can redistribute it and/or modify it. # Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import os import numpy as np diff --git a/examples/transpose/confusion_transpose/scripts/verify_data.py b/examples/transpose/scripts/verify_data.py similarity index 85% rename from examples/transpose/confusion_transpose/scripts/verify_data.py rename to examples/transpose/scripts/verify_data.py index 19122a78e348feed3cbd32480ba43320b12f6cb5..aa87fa2ce8adde5ed0492365df8afd66d4454f16 100644 --- a/examples/transpose/confusion_transpose/scripts/verify_data.py +++ b/examples/transpose/scripts/verify_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 +# This program is free software, you can redistribute it and/or modify it. # Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import sys import numpy as np diff --git a/examples/utils/init_global_memory/README.md b/examples/utils/fill/README.md similarity index 69% rename from examples/utils/init_global_memory/README.md rename to examples/utils/fill/README.md index 091646f5e2d39c24cb792eba7ae2eb61bd93ed53..bd2ccdb6393aa0ba836c47f0c3571bb589ac5080 100644 --- a/examples/utils/init_global_memory/README.md +++ b/examples/utils/fill/README.md @@ -2,13 +2,13 @@ ## 概述 -本样例介绍了调用InitGlobalMemory高阶API实现init_global_memory单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 +本样例介绍了调用Fill高阶API实现fill单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 -- 直调：使用核函数直调init_global_memory自定义算子。 +- 直调：使用核函数直调fill自定义算子。 核函数的基础调用方式，开发者完成算子核函数的开发和Tiling实现后，即可通过AscendCL运行时接口，完成算子的调用。 -- 框架调用：使用框架调用init_global_memory自定义算子。 +- 框架调用：使用框架调用fill自定义算子。 按照工程创建->算子实现->编译部署->算子调用的流程完成算子开发。整个过程都依赖于算子工程：基于工程代码框架完成算子核函数的开发和Tiling实现，通过工程编译脚本完成算子的编译部署，继而实现单算子调用或第三方框架中的算子调用。 @@ -17,7 +17,7 @@ | 调用方式 | 目录 | **描述** | | --------- | ------------------------------------------------------------ | ---------------------------------------------------------- | | 直调 | [kernel_launch_method_by_direct](./kernel_launch_method_by_direct) | host侧的核函数调用程序，包含CPU侧、NPU侧、仿真侧三种运行验证方法。 | -| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用init_global_memory算子。 | +| 框架调用 | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用fill算子。 | ## 样例支持的产品型号为： - Atlas A2训练系列产品/Atlas 800I A2推理产品 @@ -34,11 +34,11 @@ ## 算子描述 -init_global_memory单算子，用于将输出Global Memory上的数据初始化为0。两个输入作为Add计算的输入，Add的结果放入初始化为0后的Global Memory上。 +fill单算子，用于将输出Global Memory上的数据初始化为0。两个输入作为Add计算的输入，Add的结果放入初始化为0后的Global Memory上。 -init_global_memory算子规格： +fill算子规格： - + @@ -50,18 +50,18 @@ init_global_memory算子规格： - +
算子类型(OpType)InitGlobalMemoryCustom
算子类型(OpType)FillCustom
算子输入
nameshapedata typeformat
output_z256floatND
核函数名init_global_memory_custom
核函数名fill_custom
## 算子实现介绍 -本样例中实现的是固定shape为输入input_x[256]，input_y[256]，输出output_z[256]的init_global_memory算子。 +本样例中实现的是固定shape为输入input_x[256]，input_y[256]，输出output_z[256]的fill算子。 - kernel实现 - 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，首先搬运输入数据input_x、input_y到片上存储，然后使用InitGlobalMemory高阶API接口完成对Global Memory上output_z的数据初始化，并使用Add高阶API计算input_x与input_y相加的结果，将求和结果搬出到外部存储output_z上。 + 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，首先搬运输入数据input_x、input_y到片上存储，然后使用Fill高阶API接口完成对Global Memory上output_z的数据初始化，并使用Add高阶API计算input_x与input_y相加的结果，将求和结果搬出到外部存储output_z上。 - init_global_memory算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，并对Global Memory上的数据初始化，Compute任务负责对srcLocal执行add计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 + fill算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，并对Global Memory上的数据初始化，Compute任务负责对srcLocal执行add计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。 - tiling实现 diff --git a/examples/utils/init_global_memory/host_tiling/init_global_memory_custom_tiling.h b/examples/utils/fill/host_tiling/fill_custom_tiling.h similarity index 51% rename from examples/utils/init_global_memory/host_tiling/init_global_memory_custom_tiling.h rename to examples/utils/fill/host_tiling/fill_custom_tiling.h index a6f348b192c71c807b23c70744882aa25e028a18..2b72ef945c9b72a6140c33c3d8ccba2f6022aa79 100644 --- a/examples/utils/init_global_memory/host_tiling/init_global_memory_custom_tiling.h +++ b/examples/utils/fill/host_tiling/fill_custom_tiling.h @@ -1,25 +1,26 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H -#define EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H +#ifndef EXAMPLES_FILL_FILL_CUSTOM_TILING_H +#define EXAMPLES_FILL_FILL_CUSTOM_TILING_H #include "register/tilingdata_base.h" #include "tiling/tiling_api.h" #include "tiling/platform/platform_ascendc.h" namespace optiling { -BEGIN_TILING_DATA_DEF(InitGlobalMemoryCustomTilingData) +BEGIN_TILING_DATA_DEF(FillCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, k); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(InitGlobalMemoryCustom, InitGlobalMemoryCustomTilingData) +REGISTER_TILING_DATA_CLASS(FillCustom, FillCustomTilingData) } // namespace optiling -#endif // EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H +#endif // EXAMPLES_FILL_FILL_CUSTOM_TILING_H diff --git a/examples/utils/init_global_memory/kernel_impl/init_global_memory_custom.h b/examples/utils/fill/kernel_impl/fill_custom.h similarity index 70% rename from examples/utils/init_global_memory/kernel_impl/init_global_memory_custom.h rename to examples/utils/fill/kernel_impl/fill_custom.h index 21ca183c5db901e43410e57e0f15450fa03250ea..fddc7a8a773da227ef9cdb2c6349a15934626d2a 100644 --- a/examples/utils/init_global_memory/kernel_impl/init_global_memory_custom.h +++ b/examples/utils/fill/kernel_impl/fill_custom.h @@ -1,17 +1,16 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the - * "License"). Please refer to the License for details. You may not use this - * file except in compliance with the License. THIS SOFTWARE IS PROVIDED ON AN - * "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS - * FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository - * for the full text of the License. + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). + * Please refer to the License for details. You may not use this file except in compliance with the License. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_H -#define EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_H +#ifndef EXAMPLES_FILL_FILL_CUSTOM_H +#define EXAMPLES_FILL_FILL_CUSTOM_H #include "kernel_operator.h" @@ -19,18 +18,18 @@ namespace MyCustomKernel { constexpr int32_t INIT_SIZE = 256; -class KernelInitGlobalMemory { +class KernelFill { public: - __aicore__ inline KernelInitGlobalMemory() {} + __aicore__ inline KernelFill() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { xGm.SetGlobalBuffer((__gm__ float*)x + INIT_SIZE * AscendC::GetBlockIdx(), INIT_SIZE); yGm.SetGlobalBuffer((__gm__ float*)y + INIT_SIZE * AscendC::GetBlockIdx(), INIT_SIZE); zGm.SetGlobalBuffer((__gm__ float*)z + INIT_SIZE * AscendC::GetBlockIdx(), INIT_SIZE); // init zGm value - AscendC::InitGlobalMemory(zGm, INIT_SIZE, (float)(AscendC::GetBlockIdx())); + AscendC::Fill(zGm, INIT_SIZE, (float)(AscendC::GetBlockIdx())); - //需要插MTE2等MTE3的同步 + // sync of MTE2 and MTE3 is requied AscendC::TEventID eventIdMTE3ToMTE2 = GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2); AscendC::SetFlag(eventIdMTE3ToMTE2); AscendC::WaitFlag(eventIdMTE3ToMTE2); @@ -75,8 +74,8 @@ private: outQueueZ.FreeTensor(zLocal); } private: - AscendC::TQue inQueueX, inQueueY; - AscendC::TQue outQueueZ; + AscendC::TQue inQueueX, inQueueY; + AscendC::TQue outQueueZ; AscendC::GlobalTensor xGm; AscendC::GlobalTensor yGm; AscendC::GlobalTensor zGm; @@ -85,4 +84,4 @@ private: } // namespace MyCustomKernel -#endif // EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_H +#endif // EXAMPLES_FILL_FILL_CUSTOM_H diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_direct/CMakeLists.txt similarity index 82% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_direct/CMakeLists.txt index f98c46db37d9364a72146a6da000bf00d5f7fd7d..face84558d66a5b5c44056c3c9ef8adaed06a9af 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/CMakeLists.txt +++ b/examples/utils/fill/kernel_launch_method_by_direct/CMakeLists.txt @@ -18,7 +18,7 @@ if(CMAKE_INSTALL_PREFIX STREQUAL /usr/local) endif() file(GLOB KERNEL_FILES - ${CMAKE_CURRENT_SOURCE_DIR}/init_global_memory_custom.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/fill_custom.cpp ) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") @@ -30,29 +30,29 @@ else() message("invalid RUN_MODE: ${RUN_MODE}") endif() -add_executable(init_global_memory_direct_kernel_op +add_executable(fill_direct_kernel_op ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ) -target_compile_options(init_global_memory_direct_kernel_op PRIVATE +target_compile_options(fill_direct_kernel_op PRIVATE $:-g>> -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 ) -target_compile_definitions(init_global_memory_direct_kernel_op PRIVATE +target_compile_definitions(fill_direct_kernel_op PRIVATE $<$>:CUSTOM_ASCEND310P> SOC_VERSION="${SOC_VERSION}" ) -target_include_directories(init_global_memory_direct_kernel_op PRIVATE +target_include_directories(fill_direct_kernel_op PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} $:${ASCEND_CANN_PACKAGE_PATH}/include>> $:${ASCEND_CANN_PACKAGE_PATH}/runtime/include>> ) -target_link_libraries(init_global_memory_direct_kernel_op PRIVATE +target_link_libraries(fill_direct_kernel_op PRIVATE $,$>:host_intf_pub>> $:tikicpulib::${SOC_VERSION}>> $:ascendcl>> @@ -66,7 +66,7 @@ target_link_libraries(init_global_memory_direct_kernel_op PRIVATE graph_base ) -install(TARGETS init_global_memory_direct_kernel_op +install(TARGETS fill_direct_kernel_op LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/README.md b/examples/utils/fill/kernel_launch_method_by_direct/README.md similarity index 89% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/README.md rename to examples/utils/fill/kernel_launch_method_by_direct/README.md index 885c82337652545abf39054df1c5d9a7664c20f8..29b73679e95d9d4c20bbf55f51dc2382ae9e741c 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/README.md +++ b/examples/utils/fill/kernel_launch_method_by_direct/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于Kernel直调算子工程，介绍了调用InitGlobalMemory高阶API实现init_global_memory单算子，主要演示InitGlobalMemory高阶API在Kernel直调工程中的调用。 +本样例基于Kernel直调算子工程，介绍了调用Fill高阶API实现fill单算子，主要演示Fill高阶API在Kernel直调工程中的调用。 ## 目录结构介绍 | 目录及文件 | 描述 | @@ -10,7 +10,7 @@ | [cmake](./cmake) | 编译工程文件 | | [scripts](./scripts) | 包含输入数据和真值数据生成脚本文件 | | main.cpp | 主函数，调用算子的应用程序，含CPU域及NPU域调用 | -| init_global_memory_custom.cpp | 算子kernel实现 | +| fill_custom.cpp | 算子kernel实现 | | run.sh | 编译执行脚本 | | CMakeLists.txt | 编译工程文件 | diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/cmake/cpu_lib.cmake b/examples/utils/fill/kernel_launch_method_by_direct/cmake/cpu_lib.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/cmake/cpu_lib.cmake rename to examples/utils/fill/kernel_launch_method_by_direct/cmake/cpu_lib.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/cmake/npu_lib.cmake b/examples/utils/fill/kernel_launch_method_by_direct/cmake/npu_lib.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/cmake/npu_lib.cmake rename to examples/utils/fill/kernel_launch_method_by_direct/cmake/npu_lib.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/init_global_memory_custom.cpp b/examples/utils/fill/kernel_launch_method_by_direct/fill_custom.cpp similarity index 47% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/init_global_memory_custom.cpp rename to examples/utils/fill/kernel_launch_method_by_direct/fill_custom.cpp index b3261ac5b1ec5472cab603c699aa3cb44d45b0ba..40b9c12468893036f0a6816f40d2080569790e2c 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/init_global_memory_custom.cpp +++ b/examples/utils/fill/kernel_launch_method_by_direct/fill_custom.cpp @@ -1,28 +1,29 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "kernel_operator.h" -#include "../kernel_impl/init_global_memory_custom.h" +#include "../kernel_impl/fill_custom.h" -extern "C" __global__ __aicore__ void init_global_memory_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, +extern "C" __global__ __aicore__ void fill_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { - MyCustomKernel::KernelInitGlobalMemory op; + MyCustomKernel::KernelFill op; op.Init(x, y, z); op.Process(); } #ifndef ASCENDC_CPU_DEBUG // call of kernel function -void init_global_memory_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, GM_ADDR x, GM_ADDR y, GM_ADDR z, +void fill_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { - init_global_memory_custom<<>>(x, y, z, workspace, tiling); + fill_custom<<>>(x, y, z, workspace, tiling); } #endif diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/main.cpp b/examples/utils/fill/kernel_launch_method_by_direct/main.cpp similarity index 90% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/main.cpp rename to examples/utils/fill/kernel_launch_method_by_direct/main.cpp index 6997c9cf19e70629114f8b5a7a7d464334a9ea77..a0acdb3b206a097c11cb65c45c33867528a2b367 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/main.cpp +++ b/examples/utils/fill/kernel_launch_method_by_direct/main.cpp @@ -1,21 +1,22 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "../../../common/data_utils.h" #ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" -extern void init_global_memory_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, +extern void fill_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z, uint8_t *workspace, uint8_t *tiling); #else #include "tikicpulib.h" -extern "C" __global__ __aicore__ void init_global_memory_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint8_t *workspace, +extern "C" __global__ __aicore__ void fill_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint8_t *workspace, uint8_t *tiling); #endif @@ -92,7 +93,7 @@ int32_t main(int32_t argc, char *argv[]) ReadFile("../input/input_y.bin", inputSize_y, input_y, inputSize_y); AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(init_global_memory_custom, blockDim, input_x, input_y, output_z, workspace, tiling); + ICPU_RUN_KF(fill_custom, blockDim, input_x, input_y, output_z, workspace, tiling); WriteFile("../output/output_z.bin", output_z, outputSize_z); @@ -140,7 +141,7 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtMemcpy(inputXDevice, inputSize_x, inputXHost, inputSize_x, ACL_MEMCPY_HOST_TO_DEVICE)); CHECK_ACL(aclrtMemcpy(inputYDevice, inputSize_y, inputYHost, inputSize_y, ACL_MEMCPY_HOST_TO_DEVICE)); - init_global_memory_custom_do(blockDim, nullptr, stream, inputXDevice, inputYDevice, outputZDevice, + fill_custom_do(blockDim, nullptr, stream, inputXDevice, inputYDevice, outputZDevice, workspaceDevice, tilingDevice); CHECK_ACL(aclrtSynchronizeStream(stream)); diff --git a/examples/transpose/confusion_transpose/run.sh b/examples/utils/fill/kernel_launch_method_by_direct/run.sh similarity index 68% rename from examples/transpose/confusion_transpose/run.sh rename to examples/utils/fill/kernel_launch_method_by_direct/run.sh index 3116ab361999dec45d933b31b8488cbed24ac366..c8a93fc1ff0931480f8371bb7c4fe0c7e29f5124 100644 --- a/examples/transpose/confusion_transpose/run.sh +++ b/examples/utils/fill/kernel_launch_method_by_direct/run.sh @@ -1,14 +1,5 @@ #!/bin/bash -# Copyright (c) 2025 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - SHORT=r:,v:, LONG=run-mode:,soc-version:, OPTS=$(getopt -a --options $SHORT --longoptions $LONG -- "$@") @@ -35,20 +26,33 @@ rm -rf build mkdir build cd build +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== +if [ "${RUN_MODE}" = "sim" ]; then + export LD_LIBRARY_PATH=$(echo $LD_LIBRARY_PATH | sed 's/\/.*\/runtime\/lib64://g') + export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/runtime/lib64/stub:$LD_LIBRARY_PATH +fi + +source $ASCEND_HOME_DIR/bin/setenv.bash export LD_LIBRARY_PATH=${ASCEND_HOME_DIR}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH +set -euo pipefail cmake -DRUN_MODE=${RUN_MODE} -DSOC_VERSION=${SOC_VERSION} -DASCEND_CANN_PACKAGE_PATH=${ASCEND_HOME_DIR} .. make -j16 if [ "${RUN_MODE}" = "npu" ]; then - ./confusion_transpose_direct_kernel_op + ./fill_direct_kernel_op elif [ "${RUN_MODE}" = "sim" ]; then export ASCEND_TOOLKIT_HOME=${ASCEND_HOME_DIR} export ASCEND_HOME_PATH=${ASCEND_HOME_DIR} - msprof op simulator --application=./confusion_transpose_direct_kernel_op + msprof op simulator --application=./fill_direct_kernel_op elif [ "${RUN_MODE}" = "cpu" ]; then - ./confusion_transpose_direct_kernel_op + ./fill_direct_kernel_op fi - -cd .. -python3 scripts/verify_data.py output/output.bin output/golden.bin diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/gen_data.py b/examples/utils/fill/kernel_launch_method_by_direct/scripts/gen_data.py similarity index 72% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/gen_data.py rename to examples/utils/fill/kernel_launch_method_by_direct/scripts/gen_data.py index 27ffa9338e5a2478ad7f5c4c6c2ae798318966ec..77b73a313b44c1df94743b092fb16ae3ee125ac6 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/gen_data.py +++ b/examples/utils/fill/kernel_launch_method_by_direct/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 +# This program is free software, you can redistribute it and/or modify it. # Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import os import numpy as np diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_framework/CMakeLists.txt similarity index 81% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_framework/CMakeLists.txt index 584132d80993d309434fb1303de83910a1989aba..83a789d3366aa7794f09d01ba16d73c0aeac7f69 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/CMakeLists.txt +++ b/examples/utils/fill/kernel_launch_method_by_framework/CMakeLists.txt @@ -33,7 +33,14 @@ if(ENABLE_TEST AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/testcases) add_subdirectory(testcases) endif() -# modify vendor_name in install.sh and upgrade.sh +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/scripts/install.sh ${CMAKE_BINARY_DIR}/scripts/upgrade.sh COMMAND mkdir -p ${CMAKE_BINARY_DIR}/scripts COMMAND cp -r ${CMAKE_SOURCE_DIR}/scripts/* ${CMAKE_BINARY_DIR}/scripts/ diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/CMakePresets.json b/examples/utils/fill/kernel_launch_method_by_framework/CMakePresets.json similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/CMakePresets.json rename to examples/utils/fill/kernel_launch_method_by_framework/CMakePresets.json diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/README.md b/examples/utils/fill/kernel_launch_method_by_framework/README.md similarity index 94% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/README.md rename to examples/utils/fill/kernel_launch_method_by_framework/README.md index e4e186437d495789e48b066f5da66ddea64f5060..6ecf2408e979f78cddc40fdae90c82cd56b3e0a2 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/README.md +++ b/examples/utils/fill/kernel_launch_method_by_framework/README.md @@ -2,7 +2,7 @@ ## 概述 -本样例基于自定义算子工程，介绍了调用InitGlobalMemory高阶API实现init_global_memory单算子，主要演示InitGlobalMemory高阶API在自定义算子工程中的调用。 +本样例基于自定义算子工程，介绍了调用Fill高阶API实现fill单算子，主要演示Fill高阶API在自定义算子工程中的调用。 ## 目录结构 | 目录 | 描述 | @@ -74,13 +74,13 @@ cd build_out 在build_out目录下执行如下命令 ``` -./init_global_memory_custom_npu +./fill_custom_npu ``` ### 6.NPU仿真模式运行（可选） 若要执行NPU仿真，在build_out目录下执行如下命令： ``` export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH -msprof op simulator --application=./init_global_memory_custom_npu +msprof op simulator --application=./fill_custom_npu ``` 其中SOC_VERSION参数说明如下： - SOC_VERSION ：昇腾AI处理器型号，如果无法确定具体的[SOC_VERSION]，则在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询，在查询到的“Name”前增加Ascend信息，例如“Name”对应取值为xxxyy，实际配置的[SOC_VERSION]值为Ascendxxxyy。支持以下产品型号： diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/build.sh b/examples/utils/fill/kernel_launch_method_by_framework/build.sh similarity index 95% rename from examples/index/arithprogression/kernel_launch_method_by_framework/build.sh rename to examples/utils/fill/kernel_launch_method_by_framework/build.sh index 0813af07bbc6da272e38b7019e81b538bbbf8b30..229910e0133e72beadd6055c3ebb269442a6ea40 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/build.sh +++ b/examples/utils/fill/kernel_launch_method_by_framework/build.sh @@ -3,7 +3,6 @@ script_path=$(realpath $(dirname $0)) source $ASCEND_HOME_DIR/bin/setenv.bash cp -rf ../host_tiling/* op_host/ -rm -rf ./cmake/util ln -s $ASCEND_HOME_DIR/tools/op_project_templates/ascendc/customize/cmake/util/ ./cmake/util mkdir -p build_out rm -rf build_out/* @@ -70,8 +69,3 @@ else fi -# for debug -# cd build_out -# make -# cpack -# verbose append -v \ No newline at end of file diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/config.cmake b/examples/utils/fill/kernel_launch_method_by_framework/cmake/config.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/config.cmake rename to examples/utils/fill/kernel_launch_method_by_framework/cmake/config.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/func.cmake b/examples/utils/fill/kernel_launch_method_by_framework/cmake/func.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/func.cmake rename to examples/utils/fill/kernel_launch_method_by_framework/cmake/func.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/intf.cmake b/examples/utils/fill/kernel_launch_method_by_framework/cmake/intf.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/intf.cmake rename to examples/utils/fill/kernel_launch_method_by_framework/cmake/intf.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/makeself.cmake b/examples/utils/fill/kernel_launch_method_by_framework/cmake/makeself.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/cmake/makeself.cmake rename to examples/utils/fill/kernel_launch_method_by_framework/cmake/makeself.cmake diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_framework/op_host/CMakeLists.txt similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_framework/op_host/CMakeLists.txt diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom.cpp b/examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom.cpp similarity index 80% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom.cpp rename to examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom.cpp index 03c5182b4ae78e1747784e4933f05faaf857ee53..e9d29bd9496b23cc9a54d2bbd75963077cc0edd2 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom.cpp +++ b/examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom.cpp @@ -1,22 +1,23 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ #include "register/op_def_registry.h" -#include "init_global_memory_custom_tiling.h" +#include "fill_custom_tiling.h" namespace optiling { constexpr uint32_t BLOCK_DIM = 1; static ge::graphStatus TilingFunc(gert::TilingContext *context) { - InitGlobalMemoryCustomTilingData tiling; + FillCustomTilingData tiling; context->SetBlockDim(BLOCK_DIM); tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); @@ -37,9 +38,9 @@ static ge::graphStatus InferShape(gert::InferShapeContext *context) } namespace ops { -class InitGlobalMemoryCustom : public OpDef { +class FillCustom : public OpDef { public: - explicit InitGlobalMemoryCustom(const char *name) : OpDef(name) + explicit FillCustom(const char *name) : OpDef(name) { this->Input("input_x") .ParamType(REQUIRED) @@ -65,5 +66,5 @@ public: } }; -OP_ADD(InitGlobalMemoryCustom); +OP_ADD(FillCustom); } diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom_tiling.h b/examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom_tiling.h similarity index 51% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom_tiling.h rename to examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom_tiling.h index a6f348b192c71c807b23c70744882aa25e028a18..2b72ef945c9b72a6140c33c3d8ccba2f6022aa79 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_host/init_global_memory_custom_tiling.h +++ b/examples/utils/fill/kernel_launch_method_by_framework/op_host/fill_custom_tiling.h @@ -1,25 +1,26 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#ifndef EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H -#define EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H +#ifndef EXAMPLES_FILL_FILL_CUSTOM_TILING_H +#define EXAMPLES_FILL_FILL_CUSTOM_TILING_H #include "register/tilingdata_base.h" #include "tiling/tiling_api.h" #include "tiling/platform/platform_ascendc.h" namespace optiling { -BEGIN_TILING_DATA_DEF(InitGlobalMemoryCustomTilingData) +BEGIN_TILING_DATA_DEF(FillCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, k); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(InitGlobalMemoryCustom, InitGlobalMemoryCustomTilingData) +REGISTER_TILING_DATA_CLASS(FillCustom, FillCustomTilingData) } // namespace optiling -#endif // EXAMPLES_INIT_GLOBAL_MEMORY_INIT_GLOBAL_MEMORY_CUSTOM_TILING_H +#endif // EXAMPLES_FILL_FILL_CUSTOM_TILING_H diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt similarity index 80% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt index c50a409a20bd0e0cce495824295a18799e4f8be1..6f85a73857c235c8439ddfaf819164d6b9a6b952 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt +++ b/examples/utils/fill/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt @@ -1,4 +1,11 @@ -# set custom compile options +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx") add_ops_compile_options(ALL OPTIONS -g -O0) endif() diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/init_global_memory_custom.cpp b/examples/utils/fill/kernel_launch_method_by_framework/op_kernel/fill_custom.cpp similarity index 48% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/init_global_memory_custom.cpp rename to examples/utils/fill/kernel_launch_method_by_framework/op_kernel/fill_custom.cpp index 77c2fcea82d6f5de805f5b6e07130519eea02212..3f17c662606ea5657651bc81efe1c24fe33b1c5c 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/op_kernel/init_global_memory_custom.cpp +++ b/examples/utils/fill/kernel_launch_method_by_framework/op_kernel/fill_custom.cpp @@ -1,19 +1,20 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ -#include "../../../../../../kernel_impl/init_global_memory_custom.h" +#include "../../../../../../kernel_impl/fill_custom.h" -extern "C" __global__ __aicore__ void init_global_memory_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, +extern "C" __global__ __aicore__ void fill_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { - MyCustomKernel::KernelInitGlobalMemory op; + MyCustomKernel::KernelFill op; op.Init(x, y, z); op.Process(); } diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_direct/scripts/gen_data.py b/examples/utils/fill/kernel_launch_method_by_framework/scripts/gen_data.py similarity index 72% rename from examples/utils/init_global_memory/kernel_launch_method_by_direct/scripts/gen_data.py rename to examples/utils/fill/kernel_launch_method_by_framework/scripts/gen_data.py index 27ffa9338e5a2478ad7f5c4c6c2ae798318966ec..77b73a313b44c1df94743b092fb16ae3ee125ac6 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_direct/scripts/gen_data.py +++ b/examples/utils/fill/kernel_launch_method_by_framework/scripts/gen_data.py @@ -1,14 +1,14 @@ #!/usr/bin/python3 # coding=utf-8 +# This program is free software, you can redistribute it and/or modify it. # Copyright (c) 2025 Huawei Technologies Co., Ltd. # This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). # Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== import os import numpy as np diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/help.info b/examples/utils/fill/kernel_launch_method_by_framework/scripts/help.info similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/help.info rename to examples/utils/fill/kernel_launch_method_by_framework/scripts/help.info diff --git a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/install.sh b/examples/utils/fill/kernel_launch_method_by_framework/scripts/install.sh similarity index 91% rename from examples/index/arithprogression/kernel_launch_method_by_framework/scripts/install.sh rename to examples/utils/fill/kernel_launch_method_by_framework/scripts/install.sh index 8468c5a256f2c77fad5bf78ab108ca5b62aad672..48c93e67002c90b8b703c47596b6b49f2a2a006c 100644 --- a/examples/index/arithprogression/kernel_launch_method_by_framework/scripts/install.sh +++ b/examples/utils/fill/kernel_launch_method_by_framework/scripts/install.sh @@ -258,7 +258,15 @@ if [ $? -ne 0 ];then exit 1 fi -# set the set_env.bash +# This program is free software, you can redistribute it and/or modify it. +# Copyright (c) 2025 Huawei Technologies Co., Ltd. +# This file is a part of the CANN Open Software. +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. +# ====================================================================================================================== if [ -n "${INSTALL_PATH}" ] && [ -d ${INSTALL_PATH} ]; then _ASCEND_CUSTOM_OPP_PATH=${targetdir}/${vendordir} bin_path="${_ASCEND_CUSTOM_OPP_PATH}/bin" diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/upgrade.sh b/examples/utils/fill/kernel_launch_method_by_framework/scripts/upgrade.sh similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/scripts/upgrade.sh rename to examples/utils/fill/kernel_launch_method_by_framework/scripts/upgrade.sh diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_framework/testcases/CMakeLists.txt similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_framework/testcases/CMakeLists.txt diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/cmake/fun.cmake b/examples/utils/fill/kernel_launch_method_by_framework/testcases/cmake/fun.cmake similarity index 100% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/cmake/fun.cmake rename to examples/utils/fill/kernel_launch_method_by_framework/testcases/cmake/fun.cmake diff --git a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt b/examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt similarity index 64% rename from examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt rename to examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt index 6548577417a8f4bf9109c56eba49e75912e0f49c..453d17470b2df11e1d1a30d3282300bef3a4f6cd 100644 --- a/examples/select/selectwithbytesmask/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt +++ b/examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt @@ -1,10 +1,10 @@ -add_npu_target(OP SelectWithBytesMaskCustom SRC select_with_bytes_mask_custom_main.cpp) +add_npu_target(OP FillCustom SRC fill_custom_main.cpp) add_custom_target(run_npu_test COMMAND echo "===============================================================================" COMMAND echo " Run NPU test at ${CMAKE_CURRENT_BINARY_DIR}" COMMAND echo "===============================================================================" - COMMAND $ + COMMAND $ COMMAND echo "===============================================================================" ) -add_dependencies(run_npu_test select_with_bytes_mask_custom_npu) \ No newline at end of file +add_dependencies(run_npu_test fill_custom_npu) \ No newline at end of file diff --git a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/init_global_memory_custom_main.cpp b/examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/fill_custom_main.cpp similarity index 90% rename from examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/init_global_memory_custom_main.cpp rename to examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/fill_custom_main.cpp index 2eaa72329ac83f9f7cf03c98caa62ec5cc33c14f..09ded33cb23158594671ab23099505339869ee36 100644 --- a/examples/utils/init_global_memory/kernel_launch_method_by_framework/testcases/npu/init_global_memory_custom_main.cpp +++ b/examples/utils/fill/kernel_launch_method_by_framework/testcases/npu/fill_custom_main.cpp @@ -1,10 +1,11 @@ -/** +/* + * This program is free software, you can redistribute it and/or modify it. * Copyright (c) 2025 Huawei Technologies Co., Ltd. * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). + * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING + * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. * See LICENSE in the root of the software repository for the full text of the License. */ @@ -12,7 +13,7 @@ #include #include "acl/acl_rt.h" #include "acl/acl.h" -#include "aclnn_init_global_memory_custom.h" +#include "aclnn_fill_custom.h" #include "../../../../../common/data_utils.h" constexpr uint8_t SRC_NUM = 2; @@ -50,7 +51,7 @@ void DestroyStream(aclrtStream stream, int device) } struct tensorInfo { - int64_t *dims; + int64_t* dims; int64_t dimCnt; aclDataType dtype; aclFormat fmt; @@ -151,15 +152,15 @@ int main(void) size_t workspaceSize = 0; aclOpExecutor *handle; int32_t ret = 0; - ret = aclnnInitGlobalMemoryCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[INDEX_DST], + ret = aclnnFillCustomGetWorkspaceSize(tensors[0], tensors[1], tensors[INDEX_DST], &workspaceSize, &handle); - printf("aclnnInitGlobalMemoryCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); + printf("aclnnFillCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); void *workspace = nullptr; if (workspaceSize != 0) { CHECK_ACL(aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST)); } - ret = aclnnInitGlobalMemoryCustom(workspace, workspaceSize, handle, stream); - printf("aclnnInitGlobalMemoryCustom ret %u\n", ret); + ret = aclnnFillCustom(workspace, workspaceSize, handle, stream); + printf("aclnnFillCustom ret %u\n", ret); if (aclrtSynchronizeStreamWithTimeout(stream, TIMEOUT) != ACL_SUCCESS) { printf("Synchronize stream failed\n"); }