From 1ed8736e5ec61f38dc062d77532c56a733e89a66 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=A7=9C=E4=B8=9B=E9=A3=9E?= Date: Mon, 3 Mar 2025 11:02:53 +0800 Subject: [PATCH] =?UTF-8?q?=E6=95=B4=E6=94=B9=E6=A0=B7=E4=BE=8B=E5=90=8D?= =?UTF-8?q?=E7=A7=B0?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../AclNNInvocation/README.md | 10 +++--- .../AclNNInvocation/inc/common.h | 0 .../AclNNInvocation/inc/op_runner.h | 0 .../AclNNInvocation/inc/operator_desc.h | 0 .../AclNNInvocation/input/.keep | 0 .../AclNNInvocation/run.sh | 2 +- .../AclNNInvocation/scripts/acl.json | 0 .../AclNNInvocation/scripts/gen_data.py | 0 .../AclNNInvocation/scripts/verify_result.py | 0 .../AclNNInvocation/src/CMakeLists.txt | 8 ++--- .../AclNNInvocation/src/common.cpp | 0 .../AclNNInvocation/src/main.cpp | 0 .../AclNNInvocation/src/op_runner.cpp | 10 +++--- .../AclNNInvocation/src/operator_desc.cpp | 0 .../GroupBarrier.json} | 2 +- .../GroupBarrier/op_host/group_barrier.cpp} | 12 +++---- .../op_host/group_barrier_tiling.h} | 12 +++---- .../GroupBarrier/op_kernel/group_barrier.cpp} | 10 +++--- .../README.md | 34 +++++++++---------- .../install.sh | 2 +- operator/ascendc/2_features/README.md | 3 +- 21 files changed, 53 insertions(+), 52 deletions(-) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/README.md (77%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/inc/common.h (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/inc/op_runner.h (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/inc/operator_desc.h (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/input/.keep (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/run.sh (97%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/scripts/acl.json (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/scripts/gen_data.py (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/scripts/verify_result.py (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/src/CMakeLists.txt (86%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/src/common.cpp (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/src/main.cpp (100%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/src/op_runner.cpp (97%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/AclNNInvocation/src/operator_desc.cpp (100%) rename operator/ascendc/2_features/{16_cube_group_barrier/CubeGroupBarrier.json => 16_group_barrier/GroupBarrier.json} (94%) rename operator/ascendc/2_features/{16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier.cpp => 16_group_barrier/GroupBarrier/op_host/group_barrier.cpp} (87%) rename operator/ascendc/2_features/{16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier_tiling.h => 16_group_barrier/GroupBarrier/op_host/group_barrier_tiling.h} (61%) rename operator/ascendc/2_features/{16_cube_group_barrier/CubeGroupBarrier/op_kernel/cube_group_barrier.cpp => 16_group_barrier/GroupBarrier/op_kernel/group_barrier.cpp} (90%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/README.md (71%) rename operator/ascendc/2_features/{16_cube_group_barrier => 16_group_barrier}/install.sh (98%) diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/README.md b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/README.md similarity index 77% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/README.md rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/README.md index c3fe829e6..32665f595 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/README.md +++ b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/README.md @@ -1,6 +1,6 @@ ## 目录结构介绍 ``` -├── AclNNInvocation // 通过aclnn调用的方式调用CubeGroupBarrier算子 +├── AclNNInvocation // 通过aclnn调用的方式调用GroupBarrier算子 │ ├── inc // 头文件目录 │ │ ├── common.h // 声明公共方法类，用于读取二进制文件 │ │ ├── op_runner.h // 算子描述声明文件，包含算子输入/输出，算子类型以及输入描述与输出描述 @@ -25,11 +25,11 @@ 自定义算子编译部署后，会自动生成单算子API，可以直接在应用程序中调用。算子API的形式一般定义为“两段式接口”，形如： ```cpp // 获取算子使用的workspace空间大小 - aclnnStatus aclnnCubeGroupBarrierGetWorkspaceSize(const aclTensor *x, const aclTensor *y, const aclTensor *out, uint64_t workspaceSize, aclOpExecutor **executor); + aclnnStatus aclnnGroupBarrierGetWorkspaceSize(const aclTensor *x, const aclTensor *y, const aclTensor *out, uint64_t workspaceSize, aclOpExecutor **executor); // 执行算子 - aclnnStatus aclnnCubeGroupBarrier(void *workspace, int64_t workspaceSize, aclOpExecutor **executor, aclrtStream stream); + aclnnStatus aclnnGroupBarrier(void *workspace, int64_t workspaceSize, aclOpExecutor **executor, aclrtStream stream); ``` -其中aclnnCubeGroupBarrierGetWorkspaceSize为第一段接口，主要用于计算本次API调用计算过程中需要多少的workspace内存。获取到本次API计算需要的workspace大小之后，按照workspaceSize大小申请Device侧内存，然后调用第二段接口aclnnCubeGroupBarrier执行计算。具体参考[单算子API调用](https://hiascend.com/document/redirect/CannCommunityAscendCInVorkSingleOp)章节。 +其中aclnnGroupBarrierGetWorkspaceSize为第一段接口，主要用于计算本次API调用计算过程中需要多少的workspace内存。获取到本次API计算需要的workspace大小之后，按照workspaceSize大小申请Device侧内存，然后调用第二段接口aclnnGroupBarrier执行计算。具体参考[单算子API调用](https://hiascend.com/document/redirect/CannCommunityAscendCInVorkSingleOp)章节。 ## 运行样例算子 ### 1. 编译算子工程 @@ -39,7 +39,7 @@ - 进入到样例目录 以命令行方式下载样例代码，master分支为例。 ```bash - cd ${git_clone_path}/samples/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation + cd ${git_clone_path}/samples/operator/ascendc/2_features/16_group_barrier/AclNNInvocation ``` - 样例执行 diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/common.h b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/common.h similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/common.h rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/common.h diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/op_runner.h b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/op_runner.h similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/op_runner.h rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/op_runner.h diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/operator_desc.h b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/operator_desc.h similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/inc/operator_desc.h rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/inc/operator_desc.h diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/input/.keep b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/input/.keep similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/input/.keep rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/input/.keep diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/run.sh b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/run.sh similarity index 97% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/run.sh rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/run.sh index aa4b847a6..d3e54dc31 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/run.sh +++ b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/run.sh @@ -58,7 +58,7 @@ function main { echo "INFO: execute op!" check_msg="OUTPUT = 24" file_path=output_msg.txt - ./execute_cube_group_barrier | tee $file_path + ./execute_group_barrier | tee $file_path count=$(grep -c "$check_msg" $file_path) if [ $? -ne 0 ]; then diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/acl.json b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/acl.json similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/acl.json rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/acl.json diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/gen_data.py b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/gen_data.py similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/gen_data.py rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/gen_data.py diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/verify_result.py b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/verify_result.py similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/scripts/verify_result.py rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/scripts/verify_result.py diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/CMakeLists.txt b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/CMakeLists.txt similarity index 86% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/CMakeLists.txt rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/CMakeLists.txt index 79337da19..674a5e5c3 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/CMakeLists.txt +++ b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/CMakeLists.txt @@ -4,7 +4,7 @@ cmake_minimum_required(VERSION 3.5.1) # project information -project(acl_execute_cube_group_barrier) +project(acl_execute_group_barrier) # Compile options add_compile_options(-std=c++11) @@ -47,14 +47,14 @@ link_directories( ${CUST_PKG_PATH}/lib ) -add_executable(execute_cube_group_barrier +add_executable(execute_group_barrier operator_desc.cpp op_runner.cpp main.cpp common.cpp ) -target_link_libraries(execute_cube_group_barrier +target_link_libraries(execute_group_barrier ascendcl cust_opapi acl_op_compiler @@ -62,4 +62,4 @@ target_link_libraries(execute_cube_group_barrier stdc++ ) -install(TARGETS execute_cube_group_barrier DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}) +install(TARGETS execute_group_barrier DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}) diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/common.cpp b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/common.cpp similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/common.cpp rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/common.cpp diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/main.cpp b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/main.cpp similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/main.cpp rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/main.cpp diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/op_runner.cpp b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/op_runner.cpp similarity index 97% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/op_runner.cpp rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/op_runner.cpp index ec916139d..b05e2ac70 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/op_runner.cpp +++ b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/op_runner.cpp @@ -13,7 +13,7 @@ #include #include "acl/acl_op_compiler.h" -#include "aclnn_cube_group_barrier.h" +#include "aclnn_group_barrier.h" #include "common.h" using namespace std; @@ -314,13 +314,13 @@ bool OpRunner::RunOp() size_t workspaceSize = 0; aclOpExecutor *handle = nullptr; auto ret = - aclnnCubeGroupBarrierGetWorkspaceSize(inputTensor_[0], outputTensor_[0], &workspaceSize, &handle); + aclnnGroupBarrierGetWorkspaceSize(inputTensor_[0], outputTensor_[0], &workspaceSize, &handle); if (ret != ACL_SUCCESS) { (void)aclrtDestroyStream(stream); ERROR_LOG("Get Operator Workspace failed. error code is %d", static_cast(ret)); return false; } - INFO_LOG("Execute aclnnCubeGroupBarrierGetWorkspaceSize success, workspace size %lu", workspaceSize); + INFO_LOG("Execute aclnnGroupBarrierGetWorkspaceSize success, workspace size %lu", workspaceSize); if (workspaceSize != 0) { if (aclrtMalloc(&workspace_, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { @@ -328,13 +328,13 @@ bool OpRunner::RunOp() } } - ret = aclnnCubeGroupBarrier(workspace_, workspaceSize, handle, stream); + ret = aclnnGroupBarrier(workspace_, workspaceSize, handle, stream); if (ret != ACL_SUCCESS) { (void)aclrtDestroyStream(stream); ERROR_LOG("Execute Operator failed. error code is %d", static_cast(ret)); return false; } - INFO_LOG("Execute aclnnCubeGroupBarrier success"); + INFO_LOG("Execute aclnnGroupBarrier success"); ret = aclrtSynchronizeStreamWithTimeout(stream, 5000); if (ret != SUCCESS) { diff --git a/operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/operator_desc.cpp b/operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/operator_desc.cpp similarity index 100% rename from operator/ascendc/2_features/16_cube_group_barrier/AclNNInvocation/src/operator_desc.cpp rename to operator/ascendc/2_features/16_group_barrier/AclNNInvocation/src/operator_desc.cpp diff --git a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier.json b/operator/ascendc/2_features/16_group_barrier/GroupBarrier.json similarity index 94% rename from operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier.json rename to operator/ascendc/2_features/16_group_barrier/GroupBarrier.json index dbfef6c6d..284d76f28 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier.json +++ b/operator/ascendc/2_features/16_group_barrier/GroupBarrier.json @@ -1,6 +1,6 @@ [ { - "op": "CubeGroupBarrier", + "op": "GroupBarrier", "language": "cpp", "input_desc": [ { diff --git a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier.cpp b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier.cpp similarity index 87% rename from operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier.cpp rename to operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier.cpp index f429c786f..3e0bba8c2 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier.cpp +++ b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier.cpp @@ -1,5 +1,5 @@ /** - * @file cube_group_barrier.cpp + * @file group_barrier.cpp * * Copyright (C) 2023-2024. Huawei Technologies Co., Ltd. All rights reserved. * @@ -7,13 +7,13 @@ * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ -#include "cube_group_barrier_tiling.h" +#include "group_barrier_tiling.h" #include "register/op_def_registry.h" namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext *context) { - CubeGroupBarrierTilingData tiling; + GroupBarrierTilingData tiling; const gert::StorageShape* x1_shape = context->GetInputShape(0); int32_t data_sz = 1; for (int i =0; i < x1_shape->GetStorageShape().GetDimNum(); i++) @@ -38,9 +38,9 @@ static graphStatus InferShape(gert::InferShapeContext *context) } namespace ops { -class CubeGroupBarrier : public OpDef { +class GroupBarrier : public OpDef { public: - explicit CubeGroupBarrier(const char *name) : OpDef(name) + explicit GroupBarrier(const char *name) : OpDef(name) { this->Input("barworkspace") .ParamType(REQUIRED) @@ -57,5 +57,5 @@ public: .AddConfig("ascend910b"); } }; -OP_ADD(CubeGroupBarrier); +OP_ADD(GroupBarrier); } // namespace ops diff --git a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier_tiling.h b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier_tiling.h similarity index 61% rename from operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier_tiling.h rename to operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier_tiling.h index 19002a9da..94d2d40fd 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_host/cube_group_barrier_tiling.h +++ b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_host/group_barrier_tiling.h @@ -1,5 +1,5 @@ /** - * @file cube_group_barrier_tiling.h + * @file group_barrier_tiling.h * * Copyright (C) 2023-2024. Huawei Technologies Co., Ltd. All rights reserved. * @@ -7,15 +7,15 @@ * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ -#ifndef CUBE_GROUP_BARRIER_TILING_H -#define CUBE_GROUP_BARRIER_TILING_H +#ifndef GROUP_BARRIER_TILING_H +#define GROUP_BARRIER_TILING_H #include "register/tilingdata_base.h" namespace optiling { -BEGIN_TILING_DATA_DEF(CubeGroupBarrierTilingData) +BEGIN_TILING_DATA_DEF(GroupBarrierTilingData) TILING_DATA_FIELD_DEF(uint32_t, size); END_TILING_DATA_DEF; -REGISTER_TILING_DATA_CLASS(CubeGroupBarrier, CubeGroupBarrierTilingData) +REGISTER_TILING_DATA_CLASS(GroupBarrier, GroupBarrierTilingData) } // namespace optiling -#endif // CUBE_GROUP_BARRIER_TILING_H +#endif // GROUP_BARRIER_TILING_H diff --git a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_kernel/cube_group_barrier.cpp b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_kernel/group_barrier.cpp similarity index 90% rename from operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_kernel/cube_group_barrier.cpp rename to operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_kernel/group_barrier.cpp index 1a34e1eb0..2b322c73e 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/CubeGroupBarrier/op_kernel/cube_group_barrier.cpp +++ b/operator/ascendc/2_features/16_group_barrier/GroupBarrier/op_kernel/group_barrier.cpp @@ -1,5 +1,5 @@ /** - * @file cube_group_barrier.cpp + * @file group_barrier.cpp * * Copyright (C) 2022-2024. Huawei Technologies Co., Ltd. All rights reserved. * @@ -14,9 +14,9 @@ constexpr int32_t WAIT_NUM = 6; // wait num constexpr int32_t SIZE_NUM = 1; // tensor size constexpr int32_t TEMP = ARRIVE_NUM * WAIT_NUM; // initial number -class KernelCubeGroupBarrier { +class KernelGroupBarrier { public: - __aicore__ inline KernelCubeGroupBarrier() {} + __aicore__ inline KernelGroupBarrier() {} __aicore__ inline void Init(GM_ADDR barworkspace, GM_ADDR out) { this->barworkspace = barworkspace; @@ -60,10 +60,10 @@ private: GM_ADDR barworkspace; }; -extern "C" __global__ __aicore__ void cube_group_barrier(GM_ADDR barworkspace, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) +extern "C" __global__ __aicore__ void group_barrier(GM_ADDR barworkspace, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); - KernelCubeGroupBarrier op; + KernelGroupBarrier op; op.Init(barworkspace, out); op.Process(); } \ No newline at end of file diff --git a/operator/ascendc/2_features/16_cube_group_barrier/README.md b/operator/ascendc/2_features/16_group_barrier/README.md similarity index 71% rename from operator/ascendc/2_features/16_cube_group_barrier/README.md rename to operator/ascendc/2_features/16_group_barrier/README.md index 52579ad3e..36a663a62 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/README.md +++ b/operator/ascendc/2_features/16_group_barrier/README.md @@ -1,22 +1,22 @@ ## 概述 -本样例基于CubeGroupBarrier算子工程，介绍了单算子工程调用。 +本样例基于单算子工程，介绍GroupBarrier的特性。 ## 目录结构介绍 ``` -├── 16_cube_group_barrier // 使用框架调用的方式调用CubeGroupBarrier算子 -│ ├── AclNNInvocation // 通过aclnn调用的方式调用CubeGroupBarrier算子 -│ ├── CubeGroupBarrier // CubeGroupBarrier算子工程 -│ ├── CubeGroupBarrier.json // CubeGroupBarrier算子的原型定义json文件 -│ └── install.sh // 脚本，调用msOpGen生成自定义算子工程，并编译 +├── 16_group_barrier // 使用框架调用的方式调用GroupBarrier算子 +│ ├── AclNNInvocation // 通过aclnn调用的方式调用GroupBarrier算子 +│ ├── GroupBarrier // GroupBarrier算子工程 +│ ├── GroupBarrier.json // GroupBarrier算子的原型定义json文件 +│ └── install.sh // 脚本，调用msOpGen生成自定义算子工程，并编译 ``` ## 算子描述 -CubeGroupBarrier算子实现了两组存在依赖关系的AIV之间的正确同步，A组AIV计算完成后，B组AIV依赖该A组AIV的计算结果进行后续的计算，称A组AIV为Arrive组，B组AIV为Wait组。本算子不进行输入输出计算，仅通过Arrive组写完指定数值之后，Wait组读取该数值，printf打印出正确的数值完成验证。 +GroupBarrier算子实现了两组存在依赖关系的AIV之间的正确同步，A组AIV计算完成后，B组AIV依赖该A组AIV的计算结果进行后续的计算，称A组AIV为Arrive组，B组AIV为Wait组。本算子不进行输入输出计算，仅通过Arrive组写完指定数值之后，Wait组读取该数值，printf打印出正确的数值完成验证。 ## 算子规格描述 -算子CubeGroupBarrier注册的原型如下。本算子实际不进行输入输出计算，仅供参考。 +算子GroupBarrier注册的原型如下。本算子实际不进行输入输出计算，仅供参考。 - + @@ -24,7 +24,7 @@ CubeGroupBarrier算子实现了两组存在依赖关系的AIV之间的正确同 - +
算子类型(OpType)CubeGroupBarrier
算子类型(OpType)GroupBarrier
算子输入nameshapedata typeformat
x1 * 1floatND
算子输出z1 * 1floatND
核函数名cube_group_barrier
核函数名_group_barrier
## 支持的产品型号 @@ -32,15 +32,15 @@ CubeGroupBarrier算子实现了两组存在依赖关系的AIV之间的正确同 - Atlas A2训练系列产品/Atlas 800I A2推理产品 ## 算子工程介绍 -其中，算子工程目录CubeGroupBarrier包含算子的实现文件，如下所示： +其中，算子工程目录GroupBarrier包含算子的实现文件，如下所示： ``` -├── CubeGroupBarrier // CubeGroupBarrier自定义算子工程 +├── GroupBarrier // GroupBarrier自定义算子工程 │ ├── op_host // host侧实现文件 │ └── op_kernel // kernel侧实现文件 ``` -CANN软件包中提供了工程创建工具msopgen，CubeGroupBarrier算子工程可通过CubeGroupBarrier.json自动创建，自定义算子工程具体请参考[Ascend C算子开发](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)>工程化算子开发>创建算子工程 章节。 +CANN软件包中提供了工程创建工具msopgen，GroupBarrier算子工程可通过GroupBarrier.json自动创建，自定义算子工程具体请参考[Ascend C算子开发](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)>工程化算子开发>创建算子工程 章节。 -创建完自定义算子工程后，开发者重点需要完成算子host和kernel文件的功能开发。为简化样例运行流程，本样例已在CubeGroupBarrier目录中准备好了必要的算子实现，install.sh脚本会创建一个CustomOp目录，并将算子实现文件复制到对应目录下，再编译算子。 +创建完自定义算子工程后，开发者重点需要完成算子host和kernel文件的功能开发。为简化样例运行流程，本样例已在GroupBarrier目录中准备好了必要的算子实现，install.sh脚本会创建一个CustomOp目录，并将算子实现文件复制到对应目录下，再编译算子。 备注：CustomOp目录为生成目录，每次执行install.sh脚本都会删除该目录并重新生成，切勿在该目录下编码算子，会存在丢失风险。 @@ -60,7 +60,7 @@ CANN软件包中提供了工程创建工具msopgen，CubeGroupBarrier算子工 - 切换到msOpGen脚本install.sh所在目录 ```bash # 若开发者以git命令行方式clone了master分支代码，并切换目录 - cd ${git_clone_path}/samples/operator/ascendc/2_features/16_cube_group_barrier + cd ${git_clone_path}/samples/operator/ascendc/2_features/16_group_barrier ``` - 调用脚本，生成自定义算子工程，复制host和kernel实现并编译算子 @@ -114,9 +114,9 @@ CANN软件包中提供了工程创建工具msopgen，CubeGroupBarrier算子工 命令执行成功后，自定义算子包中的相关文件将部署至opp算子库环境变量ASCEND_OPP_PATH指向的的vendors/customize目录中。 ### 4. 调用执行算子工程 -- [aclnn调用CubeGroupBarrier算子工程](./AclNNInvocation/README.md) +- [aclnn调用GroupBarrier算子工程](./AclNNInvocation/README.md) ## 更新说明 | 时间 | 更新事项 | | ---------- | ---------------------------- | -| 2024/11/14 | 新增CubeGroupBarrier样例 | +| 2024/11/14 | 新增GroupBarrier样例 | | 2024/11/18 | 算子工程改写为由msOpGen生成 | diff --git a/operator/ascendc/2_features/16_cube_group_barrier/install.sh b/operator/ascendc/2_features/16_group_barrier/install.sh similarity index 98% rename from operator/ascendc/2_features/16_cube_group_barrier/install.sh rename to operator/ascendc/2_features/16_group_barrier/install.sh index 1c749dfcc..59b11009b 100644 --- a/operator/ascendc/2_features/16_cube_group_barrier/install.sh +++ b/operator/ascendc/2_features/16_group_barrier/install.sh @@ -45,7 +45,7 @@ fi source $_ASCEND_INSTALL_PATH/bin/setenv.bash export ASCEND_HOME_PATH=$_ASCEND_INSTALL_PATH -OP_NAME=CubeGroupBarrier +OP_NAME=GroupBarrier rm -rf CustomOp # Generate the op framework msopgen gen -i $OP_NAME.json -c ai_core-${SOC_VERSION} -lan cpp -out CustomOp diff --git a/operator/ascendc/2_features/README.md b/operator/ascendc/2_features/README.md index 455e14086..8c843758b 100644 --- a/operator/ascendc/2_features/README.md +++ b/operator/ascendc/2_features/README.md @@ -18,7 +18,7 @@ Ascend C相关特性的样例。特性样例逐步补充中。 | [12_cube_group](./12_cube_group) | 基于Ascend C的自定义算子及FrameworkLaunch调用样例，通过软同步控制AIC和AIV之间进行通讯，实现AI Core计算资源分组。|Atlas A2训练系列产品/Atlas 800I A2推理产品| | [13_matmul_api_ibshare](./13_matmul_api_ibshare) | 基于Ascend C的自定义Cube算子及Kernellaunch调用样例，通过A矩阵与B矩阵使能IBSHARE，实现算子性能提升|Atlas A2训练系列产品/Atlas 800I A2推理产品| | [14_matmul_api_constant](./14_matmul_api_constant) | 基于Ascend C的自定义Cube算子及FrameworkLaunch调用样例，通过使用全量常量化的MatmulApiStaticTiling模板参数，替代非常量的TCubeTiling参数，以减少Scalar计算开销，实现算子性能提升|Atlas A2训练系列产品/Atlas 800I A2推理产品| -| [16_cube_group_barrier](./16_cube_group_barrier) | 基于Ascend C的自定义Vector算子及FrameworkLaunch调用样例，通过GroupBarrier控制，实现两组AIV任务之间存在依赖关系时正确同步|Atlas A2训练系列产品/Atlas 800I A2推理产品| +| [16_group_barrier](./16_group_barrier) | 基于Ascend C的自定义Vector算子及FrameworkLaunch调用样例，通过GroupBarrier控制，实现两组AIV任务之间存在依赖关系时正确同步|Atlas A2训练系列产品/Atlas 800I A2推理产品| ## 获取样例代码 @@ -49,4 +49,5 @@ Ascend C相关特性的样例。特性样例逐步补充中。 ## 更新说明 | 时间 | 更新事项 | | ---------- | -------------------------------------------- | +| 2024/11/11 | 样例16_group_barrier名称整改 | | 2024/11/11 | 样例目录调整 | \ No newline at end of file -- Gitee