已合并
[CANNBot]Acosh适配Ascend950 #2884
Nerddddddddddd创建于 5月21日
[CANNBot]Acosh适配Ascend950 #2884
已合并
共 21 个文件变更+1482-51
| @@ -1,19 +1,18 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | 1 | +# --------------------------------------------------------------------------------------------------------- |
| 2 | -# This program is free software, you can redistribute it and/or modify it. | 2 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# This file is a part of the CANN Open Software. | 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 7 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | # See LICENSE in the root of the software repository for the full text of the License. | 8 | # See LICENSE in the root of the software repository for the full text of the License. |
| 10 | -# ---------------------------------------------------------------------------- | 9 | +# --------------------------------------------------------------------------------------------------------- |
| 11 | -file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | 10 | + |
| 12 | -foreach(SUB_DIR ${CURRENT_DIRS}) | 11 | +# Generated By CANNBot |
| 13 | - if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | 12 | + |
| 14 | - add_subdirectory(${SUB_DIR}) | 13 | +set(SUPPORT_COMPUTE_UNIT "ascend950") |
| 15 | - endif() | 14 | +set(SUPPORT_TILING_DIR "arch35") |
| 16 | -endforeach() | 15 | +add_all_modules_sources(OPTYPE acosh ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE) |
| 17 | 16 | ||
| 18 | if(ENABLE_TEST AND (UT_TEST_ALL OR OP_KERNEL_AICPU_UT)) | 17 | if(ENABLE_TEST AND (UT_TEST_ALL OR OP_KERNEL_AICPU_UT)) |
| 19 | list(FIND ASCEND_OP_NAME acosh ACOSH_INDEX) | 18 | list(FIND ASCEND_OP_NAME acosh ACOSH_INDEX) |
| @@ -1,3 +1,93 @@ | |||
| 1 | # Acosh | 1 | # Acosh |
| 2 | 2 | ||
| 3 | -本目录仅包含Acosh算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。 | 3 | +> 反双曲余弦（Inverse Hyperbolic Cosine）算子，对应 PyTorch `torch.acosh` / `torch.acosh_`。 |
| 4 | +> 在线 aclnn API 文档：[docs/aclnnAcosh&aclnnInplaceAcosh.md](docs/aclnnAcosh&aclnnInplaceAcosh.md) | ||
| 5 | + | ||
| 6 | +## 产品支持情况 | ||
| 7 | + | ||
| 8 | +|产品 | 是否支持 | | ||
| 9 | +|:-------------------------|:----------:| | ||
| 10 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | | ||
| 11 | +| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | | ||
| 12 | +| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | | ||
| 13 | +| <term>Atlas 200I/500 A2 推理产品</term> | × | | ||
| 14 | +| <term>Atlas 推理系列产品</term> | √ | | ||
| 15 | +| <term>Atlas 训练系列产品</term> | √ | | ||
| 16 | + | ||
| 17 | +## 功能说明 | ||
| 18 | + | ||
| 19 | +- 接口功能：对输入Tensor中的每个元素进行反双曲余弦操作后输出。 | ||
| 20 | + | ||
| 21 | +- 计算公式： | ||
| 22 | + | ||
| 23 | + $$ | ||
| 24 | + out=cosh^{-1}(self) | ||
| 25 | + $$ | ||
| 26 | + | ||
| 27 | +## 参数说明 | ||
| 28 | + | ||
| 29 | +<table style="table-layout: fixed; width: 1000px"><colgroup> | ||
| 30 | +<col style="width: 50px"> | ||
| 31 | +<col style="width: 70px"> | ||
| 32 | +<col style="width: 200px"> | ||
| 33 | +<col style="width: 100px"> | ||
| 34 | +<col style="width: 50px"> | ||
| 35 | +</colgroup> | ||
| 36 | +<thead> | ||
| 37 | + <tr> | ||
| 38 | + <th>参数名</th> | ||
| 39 | + <th>输入/输出/属性</th> | ||
| 40 | + <th>描述</th> | ||
| 41 | + <th>数据类型</th> | ||
| 42 | + <th>数据格式</th> | ||
| 43 | + </tr></thead> | ||
| 44 | +<tbody> | ||
| 45 | + <tr> | ||
| 46 | + <td>self</td> | ||
| 47 | + <td>输入</td> | ||
| 48 | + <td>公式中的输入 self，dtype 与 out 一致，shape 任意（0~8 维），支持空 Tensor，支持非连续 Tensor。</td> | ||
| 49 | + <td>FLOAT、FLOAT16、BFLOAT16</td> | ||
| 50 | + <td>ND</td> | ||
| 51 | + </tr> | ||
| 52 | + <tr> | ||
| 53 | + <td>out</td> | ||
| 54 | + <td>输出</td> | ||
| 55 | + <td>公式中的输出 out，shape 与 self 完全一致；非原地接口专有。</td> | ||
| 56 | + <td>FLOAT、FLOAT16、BFLOAT16</td> | ||
| 57 | + <td>ND</td> | ||
| 58 | + </tr> | ||
| 59 | + <tr> | ||
| 60 | + <td>selfRef</td> | ||
| 61 | + <td>输入/输出</td> | ||
| 62 | + <td>原地接口专有，既作输入又作输出的张量，计算结果原地写回。</td> | ||
| 63 | + <td>FLOAT、FLOAT16、BFLOAT16</td> | ||
| 64 | + <td>ND</td> | ||
| 65 | + </tr> | ||
| 66 | +</tbody></table> | ||
| 67 | + | ||
| 68 | +- 参数 `self` 与 `out` 的 dtype 必须完全一致；shape 维度数 ∈ [0, 8]。 | ||
| 69 | +- 仅支持 ND 数据格式；不支持 NCHW / NHWC / 5HD 等私有格式。 | ||
| 70 | + | ||
| 71 | +## 约束说明 | ||
| 72 | + | ||
| 73 | + | ||
| 74 | +## 调用说明 | ||
| 75 | + | ||
| 76 | +<table><thead> | ||
| 77 | + <tr> | ||
| 78 | + <th>调用方式</th> | ||
| 79 | + <th>调用样例</th> | ||
| 80 | + <th>说明</th> | ||
| 81 | + </tr></thead> | ||
| 82 | +<tbody> | ||
| 83 | + <tr> | ||
| 84 | + <td>aclnn 调用（非原地）</td> | ||
| 85 | + <td><a href="./examples/test_aclnn_acosh.cpp">examples/test_aclnn_acosh.cpp</a></td> | ||
| 86 | + <td rowspan="2">参考 ops-math 仓库根 build.sh 完成算子编译与算子包安装，再用 examples 验证（NPU 真机）。</td> | ||
| 87 | + </tr> | ||
| 88 | + <tr> | ||
| 89 | + <td>aclnn 调用（原地）</td> | ||
| 90 | + <td><a href="./examples/test_aclnn_inplace_acosh.cpp">examples/test_aclnn_inplace_acosh.cpp</a></td> | ||
| 91 | + </tr> | ||
| 92 | +</tbody></table> | ||
| 93 | + | ||
| @@ -4,14 +4,14 @@ | |||
| 4 | 4 | ||
| 5 | ## 产品支持情况 | 5 | ## 产品支持情况 |
| 6 | 6 | ||
| 7 | -| 产品 | 是否支持 | | 7 | +|产品 | 是否支持 | |
| 8 | -| :----------------------------------------------------------- | :------: | | 8 | +|:-------------------------|:----------:| |
| 9 | -| <term>Ascend 950PR/Ascend 950DT</term> | × | | 9 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | |
| 10 | -| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | | 10 | +| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | |
| 11 | -| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | | 11 | +| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | |
| 12 | -| <term>Atlas 200I/500 A2 推理产品</term> | × | | 12 | +| <term>Atlas 200I/500 A2 推理产品</term> | × | |
| 13 | -| <term>Atlas 推理系列产品</term> | √ | | 13 | +| <term>Atlas 推理系列产品</term> | √ | |
| 14 | -| <term>Atlas 训练系列产品</term> | √ | | 14 | +| <term>Atlas 训练系列产品</term> | √ | |
| 15 | 15 | ||
| 16 | ## 功能说明 | 16 | ## 功能说明 |
| 17 | 17 | ||
| @@ -32,33 +32,33 @@ | |||
| 32 | 32 | ||
| 33 | ```cpp | 33 | ```cpp |
| 34 | aclnnStatus aclnnAcoshGetWorkspaceSize( | 34 | aclnnStatus aclnnAcoshGetWorkspaceSize( |
| 35 | - const aclTensor* self, | 35 | + const aclTensor *self, |
| 36 | - aclTensor* out, | 36 | + aclTensor *out, |
| 37 | - uint64_t* workspaceSize, | 37 | + uint64_t *workspaceSize, |
| 38 | - aclOpExecutor** executor) | 38 | + aclOpExecutor **executor) |
| 39 | ``` | 39 | ``` |
| 40 | 40 | ||
| 41 | ```cpp | 41 | ```cpp |
| 42 | aclnnStatus aclnnAcosh( | 42 | aclnnStatus aclnnAcosh( |
| 43 | - void* workspace, | 43 | + void *workspace, |
| 44 | - uint64_t workspaceSize, | 44 | + uint64_t workspaceSize, |
| 45 | - aclOpExecutor* executor, | 45 | + aclOpExecutor *executor, |
| 46 | - aclrtStream stream) | 46 | + aclrtStream stream) |
| 47 | ``` | 47 | ``` |
| 48 | 48 | ||
| 49 | ```cpp | 49 | ```cpp |
| 50 | aclnnStatus aclnnInplaceAcoshGetWorkspaceSize( | 50 | aclnnStatus aclnnInplaceAcoshGetWorkspaceSize( |
| 51 | - aclTensor* selfRef, | 51 | + aclTensor *selfRef, |
| 52 | - uint64_t* workspaceSize, | 52 | + uint64_t *workspaceSize, |
| 53 | - aclOpExecutor** executor) | 53 | + aclOpExecutor **executor) |
| 54 | ``` | 54 | ``` |
| 55 | 55 | ||
| 56 | ```cpp | 56 | ```cpp |
| 57 | aclnnStatus aclnnInplaceAcosh( | 57 | aclnnStatus aclnnInplaceAcosh( |
| 58 | - void* workspace, | 58 | + void *workspace, |
| 59 | - uint64_t workspaceSize, | 59 | + uint64_t workspaceSize, |
| 60 | - aclOpExecutor* executor, | 60 | + aclOpExecutor *executor, |
| 61 | - aclrtStream stream) | 61 | + aclrtStream stream) |
| 62 | ``` | 62 | ``` |
| 63 | 63 | ||
| 64 | ## aclnnAcoshGetWorkspaceSize | 64 | ## aclnnAcoshGetWorkspaceSize |
| @@ -347,6 +347,7 @@ aclnnStatus aclnnInplaceAcosh( | |||
| 347 | 347 | ||
| 348 | aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。 | 348 | aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。 |
| 349 | 349 | ||
| 350 | + | ||
| 350 | ## 约束说明 | 351 | ## 约束说明 |
| 351 | 352 | ||
| 352 | - 确定性计算： | 353 | - 确定性计算： |
| @@ -513,3 +514,4 @@ int main() { | |||
| 513 | return 0; | 514 | return 0; |
| 514 | } | 515 | } |
| 515 | ``` | 516 | ``` |
| 517 | + | ||
| @@ -0,0 +1,182 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/** | ||
| 14 | + * @file test_aclnn_acosh.cpp | ||
| 15 | + * @brief aclnnAcosh（非原地）调用示例 | ||
| 16 | + * | ||
| 17 | + * 演示标准 aclnn 两段式调用流程： | ||
| 18 | + * 1) aclInit / aclrtSetDevice / aclrtCreateStream | ||
| 19 | + * 2) Host 构造输入张量并拷贝到 Device | ||
| 20 | + * 3) aclnnAcoshGetWorkspaceSize → workspaceSize / executor | ||
| 21 | + * 4) 申请 workspace 后调用 aclnnAcosh | ||
| 22 | + * 5) aclrtSynchronizeStream 等待执行结束 | ||
| 23 | + * 6) Device → Host 拷回输出并打印 | ||
| 24 | + * 7) 资源释放 + aclFinalize | ||
| 25 | + * | ||
| 26 | + * 前置条件：算子已编译并安装到 ${ASCEND_HOME_PATH}/opp/vendors/acosh_custom/ | ||
| 27 | + */ | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + do { \ | ||
| 39 | + if (!(cond)) { \ | ||
| 40 | + return_expr; \ | ||
| 41 | + } \ | ||
| 42 | + } while (0) | ||
| 43 | + | ||
| 44 | + | ||
| 45 | + do { \ | ||
| 46 | + printf(message, ##__VA_ARGS__); \ | ||
| 47 | + } while (0) | ||
| 48 | + | ||
| 49 | +static int64_t GetShapeSize(const std::vector<int64_t> &shape) | ||
| 50 | +{ | ||
| 51 | + int64_t size = 1; | ||
| 52 | + for (auto d : shape) { | ||
| 53 | + size *= d; | ||
| 54 | + } | ||
| 55 | + return size; | ||
| 56 | +} | ||
| 57 | + | ||
| 58 | +static int Init(int32_t deviceId, aclrtStream *stream) | ||
| 59 | +{ | ||
| 60 | + auto ret = aclInit(nullptr); | ||
| 61 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 62 | + ret = aclrtSetDevice(deviceId); | ||
| 63 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 64 | + ret = aclrtCreateStream(stream); | ||
| 65 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 66 | + return 0; | ||
| 67 | +} | ||
| 68 | + | ||
| 69 | +template <typename T> | ||
| 70 | +static int CreateAclTensor(const std::vector<T> &hostData, | ||
| 71 | + const std::vector<int64_t> &shape, | ||
| 72 | + void **deviceAddr, | ||
| 73 | + aclDataType dataType, | ||
| 74 | + aclTensor **tensor) | ||
| 75 | +{ | ||
| 76 | + auto bytes = GetShapeSize(shape) * sizeof(T); | ||
| 77 | + auto ret = aclrtMalloc(deviceAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 78 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 79 | + ret = aclrtMemcpy(*deviceAddr, bytes, hostData.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 80 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy H2D failed. ERROR: %d\n", ret); return ret); | ||
| 81 | + | ||
| 82 | + // 行主序连续 strides | ||
| 83 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 84 | + for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) { | ||
| 85 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 86 | + } | ||
| 87 | + | ||
| 88 | + *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, | ||
| 89 | + strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 90 | + shape.data(), shape.size(), *deviceAddr); | ||
| 91 | + return 0; | ||
| 92 | +} | ||
| 93 | + | ||
| 94 | +int main() | ||
| 95 | +{ | ||
| 96 | + // 1. 初始化 ACL | ||
| 97 | + int32_t deviceId = 0; | ||
| 98 | + aclrtStream stream = nullptr; | ||
| 99 | + auto ret = Init(deviceId, &stream); | ||
| 100 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 101 | + | ||
| 102 | + // 2. 构造输入：shape = [2, 4]，全部位于定义域 [1, +∞) | ||
| 103 | + std::vector<int64_t> selfShape = {2, 4}; | ||
| 104 | + std::vector<int64_t> outShape = {2, 4}; | ||
| 105 | + std::vector<float> selfHostData = {1.0f, 2.0f, 3.0f, 5.0f, | ||
| 106 | + 10.0f, 1.5f, 2.5f, 4.0f}; | ||
| 107 | + std::vector<float> outHostData(static_cast<size_t>(GetShapeSize(outShape)), 0.0f); | ||
| 108 | + | ||
| 109 | + aclTensor *selfTensor = nullptr; | ||
| 110 | + void *selfDevice = nullptr; | ||
| 111 | + ret = CreateAclTensor(selfHostData, selfShape, &selfDevice, | ||
| 112 | + aclDataType::ACL_FLOAT, &selfTensor); | ||
| 113 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 114 | + | ||
| 115 | + aclTensor *outTensor = nullptr; | ||
| 116 | + void *outDevice = nullptr; | ||
| 117 | + ret = CreateAclTensor(outHostData, outShape, &outDevice, | ||
| 118 | + aclDataType::ACL_FLOAT, &outTensor); | ||
| 119 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 120 | + | ||
| 121 | + // 3. 第一段接口：获取 workspace 大小 + executor | ||
| 122 | + uint64_t workspaceSize = 0; | ||
| 123 | + aclOpExecutor *executor = nullptr; | ||
| 124 | + ret = aclnnAcoshGetWorkspaceSize(selfTensor, outTensor, &workspaceSize, &executor); | ||
| 125 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 126 | + LOG_PRINT("aclnnAcoshGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 127 | + | ||
| 128 | + // 4. 申请 workspace（如有） | ||
| 129 | + void *workspaceAddr = nullptr; | ||
| 130 | + if (workspaceSize > 0) { | ||
| 131 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 132 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 133 | + LOG_PRINT("workspace aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 134 | + } | ||
| 135 | + | ||
| 136 | + // 5. 第二段接口：实际执行算子 | ||
| 137 | + ret = aclnnAcosh(workspaceAddr, workspaceSize, executor, stream); | ||
| 138 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 139 | + LOG_PRINT("aclnnAcosh failed. ERROR: %d\n", ret); return ret); | ||
| 140 | + | ||
| 141 | + // 6. 同步等待 | ||
| 142 | + ret = aclrtSynchronizeStream(stream); | ||
| 143 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 144 | + LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 145 | + | ||
| 146 | + // 7. 拷回输出并打印（同时用 std::acosh 粗略对照） | ||
| 147 | + auto outBytes = GetShapeSize(outShape) * sizeof(float); | ||
| 148 | + std::vector<float> npuResult(static_cast<size_t>(GetShapeSize(outShape)), 0.0f); | ||
| 149 | + ret = aclrtMemcpy(npuResult.data(), outBytes, outDevice, outBytes, | ||
| 150 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 151 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 152 | + LOG_PRINT("aclrtMemcpy D2H failed. ERROR: %d\n", ret); return ret); | ||
| 153 | + | ||
| 154 | + LOG_PRINT("\n=== aclnnAcosh result (shape=[2,4], dtype=FLOAT) ===\n"); | ||
| 155 | + int allClose = 1; | ||
| 156 | + for (size_t i = 0; i < npuResult.size(); ++i) { | ||
| 157 | + float expected = std::acosh(selfHostData[i]); | ||
| 158 | + float diff = std::fabs(npuResult[i] - expected); | ||
| 159 | + // FP32 阈值参考：MERE 2^-13 ≈ 1.22e-4，这里取 1e-4 作为示例粗对照 | ||
| 160 | + const char *flag = (diff < 1e-4f) ? "OK" : "DIFF"; | ||
| 161 | + if (diff >= 1e-4f) { | ||
| 162 | + allClose = 0; | ||
| 163 | + } | ||
| 164 | + LOG_PRINT(" self[%zu]=%-9.4f npu=%-12.7f std::acosh=%-12.7f diff=%-10.3e [%s]\n", | ||
| 165 | + i, selfHostData[i], npuResult[i], expected, diff, flag); | ||
| 166 | + } | ||
| 167 | + LOG_PRINT("=== aclnnAcosh %s ===\n", allClose ? "PASS (粗对照)" : "DIFF (粗对照)"); | ||
| 168 | + | ||
| 169 | + // 8. 释放资源 | ||
| 170 | + aclDestroyTensor(selfTensor); | ||
| 171 | + aclDestroyTensor(outTensor); | ||
| 172 | + aclrtFree(selfDevice); | ||
| 173 | + aclrtFree(outDevice); | ||
| 174 | + if (workspaceSize > 0 && workspaceAddr != nullptr) { | ||
| 175 | + aclrtFree(workspaceAddr); | ||
| 176 | + } | ||
| 177 | + aclrtDestroyStream(stream); | ||
| 178 | + aclrtResetDevice(deviceId); | ||
| 179 | + aclFinalize(); | ||
| 180 | + | ||
| 181 | + return allClose ? 0 : 1; | ||
| 182 | +} | ||
| @@ -0,0 +1,167 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/** | ||
| 14 | + * @file test_aclnn_inplace_acosh.cpp | ||
| 15 | + * @brief aclnnInplaceAcosh（原地）调用示例 | ||
| 16 | + * | ||
| 17 | + * 与 test_aclnn_acosh.cpp 的差异： | ||
| 18 | + * - 只有 selfRef 一个张量，既作输入又作输出，结果原地写回； | ||
| 19 | + * - 调用 aclnnInplaceAcoshGetWorkspaceSize / aclnnInplaceAcosh。 | ||
| 20 | + * | ||
| 21 | + * 前置条件：算子已编译并安装到 ${ASCEND_HOME_PATH}/opp/vendors/acosh_custom/ | ||
| 22 | + */ | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + do { \ | ||
| 34 | + if (!(cond)) { \ | ||
| 35 | + return_expr; \ | ||
| 36 | + } \ | ||
| 37 | + } while (0) | ||
| 38 | + | ||
| 39 | + | ||
| 40 | + do { \ | ||
| 41 | + printf(message, ##__VA_ARGS__); \ | ||
| 42 | + } while (0) | ||
| 43 | + | ||
| 44 | +static int64_t GetShapeSize(const std::vector<int64_t> &shape) | ||
| 45 | +{ | ||
| 46 | + int64_t size = 1; | ||
| 47 | + for (auto d : shape) { | ||
| 48 | + size *= d; | ||
| 49 | + } | ||
| 50 | + return size; | ||
| 51 | +} | ||
| 52 | + | ||
| 53 | +static int Init(int32_t deviceId, aclrtStream *stream) | ||
| 54 | +{ | ||
| 55 | + auto ret = aclInit(nullptr); | ||
| 56 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 57 | + ret = aclrtSetDevice(deviceId); | ||
| 58 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 59 | + ret = aclrtCreateStream(stream); | ||
| 60 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 61 | + return 0; | ||
| 62 | +} | ||
| 63 | + | ||
| 64 | +template <typename T> | ||
| 65 | +static int CreateAclTensor(const std::vector<T> &hostData, | ||
| 66 | + const std::vector<int64_t> &shape, | ||
| 67 | + void **deviceAddr, | ||
| 68 | + aclDataType dataType, | ||
| 69 | + aclTensor **tensor) | ||
| 70 | +{ | ||
| 71 | + auto bytes = GetShapeSize(shape) * sizeof(T); | ||
| 72 | + auto ret = aclrtMalloc(deviceAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 73 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 74 | + ret = aclrtMemcpy(*deviceAddr, bytes, hostData.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 75 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy H2D failed. ERROR: %d\n", ret); return ret); | ||
| 76 | + | ||
| 77 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 78 | + for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) { | ||
| 79 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 80 | + } | ||
| 81 | + | ||
| 82 | + *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, | ||
| 83 | + strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 84 | + shape.data(), shape.size(), *deviceAddr); | ||
| 85 | + return 0; | ||
| 86 | +} | ||
| 87 | + | ||
| 88 | +int main() | ||
| 89 | +{ | ||
| 90 | + // 1. 初始化 ACL | ||
| 91 | + int32_t deviceId = 0; | ||
| 92 | + aclrtStream stream = nullptr; | ||
| 93 | + auto ret = Init(deviceId, &stream); | ||
| 94 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 95 | + | ||
| 96 | + // 2. 构造 selfRef：shape=[2,4]，同时作为输入与输出 | ||
| 97 | + std::vector<int64_t> selfShape = {2, 4}; | ||
| 98 | + std::vector<float> selfHost = {1.0f, 2.0f, 3.0f, 5.0f, | ||
| 99 | + 10.0f, 1.5f, 2.5f, 4.0f}; | ||
| 100 | + // 保留原始值用于后续 std::acosh 粗对照 | ||
| 101 | + const std::vector<float> origSelf = selfHost; | ||
| 102 | + | ||
| 103 | + aclTensor *selfRef = nullptr; | ||
| 104 | + void *selfDevice = nullptr; | ||
| 105 | + ret = CreateAclTensor(selfHost, selfShape, &selfDevice, | ||
| 106 | + aclDataType::ACL_FLOAT, &selfRef); | ||
| 107 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 108 | + | ||
| 109 | + // 3. 第一段接口 | ||
| 110 | + uint64_t workspaceSize = 0; | ||
| 111 | + aclOpExecutor *executor = nullptr; | ||
| 112 | + ret = aclnnInplaceAcoshGetWorkspaceSize(selfRef, &workspaceSize, &executor); | ||
| 113 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 114 | + LOG_PRINT("aclnnInplaceAcoshGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 115 | + | ||
| 116 | + // 4. 申请 workspace | ||
| 117 | + void *workspaceAddr = nullptr; | ||
| 118 | + if (workspaceSize > 0) { | ||
| 119 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 120 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 121 | + LOG_PRINT("workspace aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 122 | + } | ||
| 123 | + | ||
| 124 | + // 5. 第二段接口：原地执行 | ||
| 125 | + ret = aclnnInplaceAcosh(workspaceAddr, workspaceSize, executor, stream); | ||
| 126 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 127 | + LOG_PRINT("aclnnInplaceAcosh failed. ERROR: %d\n", ret); return ret); | ||
| 128 | + | ||
| 129 | + // 6. 同步等待 | ||
| 130 | + ret = aclrtSynchronizeStream(stream); | ||
| 131 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 132 | + LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 133 | + | ||
| 134 | + // 7. 拷回（注意：selfRef 内存即输出） | ||
| 135 | + auto bytes = GetShapeSize(selfShape) * sizeof(float); | ||
| 136 | + std::vector<float> npuResult(static_cast<size_t>(GetShapeSize(selfShape)), 0.0f); | ||
| 137 | + ret = aclrtMemcpy(npuResult.data(), bytes, selfDevice, bytes, | ||
| 138 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 139 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 140 | + LOG_PRINT("aclrtMemcpy D2H failed. ERROR: %d\n", ret); return ret); | ||
| 141 | + | ||
| 142 | + LOG_PRINT("\n=== aclnnInplaceAcosh result (shape=[2,4], dtype=FLOAT) ===\n"); | ||
| 143 | + int allClose = 1; | ||
| 144 | + for (size_t i = 0; i < npuResult.size(); ++i) { | ||
| 145 | + float expected = std::acosh(origSelf[i]); | ||
| 146 | + float diff = std::fabs(npuResult[i] - expected); | ||
| 147 | + const char *flag = (diff < 1e-4f) ? "OK" : "DIFF"; | ||
| 148 | + if (diff >= 1e-4f) { | ||
| 149 | + allClose = 0; | ||
| 150 | + } | ||
| 151 | + LOG_PRINT(" selfRef[%zu] in=%-9.4f out(原地)=%-12.7f std::acosh=%-12.7f diff=%-10.3e [%s]\n", | ||
| 152 | + i, origSelf[i], npuResult[i], expected, diff, flag); | ||
| 153 | + } | ||
| 154 | + LOG_PRINT("=== aclnnInplaceAcosh %s ===\n", allClose ? "PASS (粗对照)" : "DIFF (粗对照)"); | ||
| 155 | + | ||
| 156 | + // 8. 释放资源 | ||
| 157 | + aclDestroyTensor(selfRef); | ||
| 158 | + aclrtFree(selfDevice); | ||
| 159 | + if (workspaceSize > 0 && workspaceAddr != nullptr) { | ||
| 160 | + aclrtFree(workspaceAddr); | ||
| 161 | + } | ||
| 162 | + aclrtDestroyStream(stream); | ||
| 163 | + aclrtResetDevice(deviceId); | ||
| 164 | + aclFinalize(); | ||
| 165 | + | ||
| 166 | + return allClose ? 0 : 1; | ||
| 167 | +} | ||
| @@ -0,0 +1,182 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/** | ||
| 14 | + * @file test_aclnn_acosh.cpp | ||
| 15 | + * @brief aclnnAcosh（非原地）调用示例 | ||
| 16 | + * | ||
| 17 | + * 演示标准 aclnn 两段式调用流程： | ||
| 18 | + * 1) aclInit / aclrtSetDevice / aclrtCreateStream | ||
| 19 | + * 2) Host 构造输入张量并拷贝到 Device | ||
| 20 | + * 3) aclnnAcoshGetWorkspaceSize → workspaceSize / executor | ||
| 21 | + * 4) 申请 workspace 后调用 aclnnAcosh | ||
| 22 | + * 5) aclrtSynchronizeStream 等待执行结束 | ||
| 23 | + * 6) Device → Host 拷回输出并打印 | ||
| 24 | + * 7) 资源释放 + aclFinalize | ||
| 25 | + * | ||
| 26 | + * 前置条件：算子已编译并安装到 ${ASCEND_HOME_PATH}/opp/vendors/acosh_custom/ | ||
| 27 | + */ | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + do { \ | ||
| 39 | + if (!(cond)) { \ | ||
| 40 | + return_expr; \ | ||
| 41 | + } \ | ||
| 42 | + } while (0) | ||
| 43 | + | ||
| 44 | + | ||
| 45 | + do { \ | ||
| 46 | + printf(message, ##__VA_ARGS__); \ | ||
| 47 | + } while (0) | ||
| 48 | + | ||
| 49 | +static int64_t GetShapeSize(const std::vector<int64_t> &shape) | ||
| 50 | +{ | ||
| 51 | + int64_t size = 1; | ||
| 52 | + for (auto d : shape) { | ||
| 53 | + size *= d; | ||
| 54 | + } | ||
| 55 | + return size; | ||
| 56 | +} | ||
| 57 | + | ||
| 58 | +static int Init(int32_t deviceId, aclrtStream *stream) | ||
| 59 | +{ | ||
| 60 | + auto ret = aclInit(nullptr); | ||
| 61 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 62 | + ret = aclrtSetDevice(deviceId); | ||
| 63 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 64 | + ret = aclrtCreateStream(stream); | ||
| 65 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 66 | + return 0; | ||
| 67 | +} | ||
| 68 | + | ||
| 69 | +template <typename T> | ||
| 70 | +static int CreateAclTensor(const std::vector<T> &hostData, | ||
| 71 | + const std::vector<int64_t> &shape, | ||
| 72 | + void **deviceAddr, | ||
| 73 | + aclDataType dataType, | ||
| 74 | + aclTensor **tensor) | ||
| 75 | +{ | ||
| 76 | + auto bytes = GetShapeSize(shape) * sizeof(T); | ||
| 77 | + auto ret = aclrtMalloc(deviceAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 78 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 79 | + ret = aclrtMemcpy(*deviceAddr, bytes, hostData.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 80 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy H2D failed. ERROR: %d\n", ret); return ret); | ||
| 81 | + | ||
| 82 | + // 行主序连续 strides | ||
| 83 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 84 | + for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) { | ||
| 85 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 86 | + } | ||
| 87 | + | ||
| 88 | + *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, | ||
| 89 | + strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 90 | + shape.data(), shape.size(), *deviceAddr); | ||
| 91 | + return 0; | ||
| 92 | +} | ||
| 93 | + | ||
| 94 | +int main() | ||
| 95 | +{ | ||
| 96 | + // 1. 初始化 ACL | ||
| 97 | + int32_t deviceId = 0; | ||
| 98 | + aclrtStream stream = nullptr; | ||
| 99 | + auto ret = Init(deviceId, &stream); | ||
| 100 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 101 | + | ||
| 102 | + // 2. 构造输入：shape = [2, 4]，全部位于定义域 [1, +∞) | ||
| 103 | + std::vector<int64_t> selfShape = {2, 4}; | ||
| 104 | + std::vector<int64_t> outShape = {2, 4}; | ||
| 105 | + std::vector<float> selfHostData = {1.0f, 2.0f, 3.0f, 5.0f, | ||
| 106 | + 10.0f, 1.5f, 2.5f, 4.0f}; | ||
| 107 | + std::vector<float> outHostData(static_cast<size_t>(GetShapeSize(outShape)), 0.0f); | ||
| 108 | + | ||
| 109 | + aclTensor *selfTensor = nullptr; | ||
| 110 | + void *selfDevice = nullptr; | ||
| 111 | + ret = CreateAclTensor(selfHostData, selfShape, &selfDevice, | ||
| 112 | + aclDataType::ACL_FLOAT, &selfTensor); | ||
| 113 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 114 | + | ||
| 115 | + aclTensor *outTensor = nullptr; | ||
| 116 | + void *outDevice = nullptr; | ||
| 117 | + ret = CreateAclTensor(outHostData, outShape, &outDevice, | ||
| 118 | + aclDataType::ACL_FLOAT, &outTensor); | ||
| 119 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 120 | + | ||
| 121 | + // 3. 第一段接口：获取 workspace 大小 + executor | ||
| 122 | + uint64_t workspaceSize = 0; | ||
| 123 | + aclOpExecutor *executor = nullptr; | ||
| 124 | + ret = aclnnAcoshGetWorkspaceSize(selfTensor, outTensor, &workspaceSize, &executor); | ||
| 125 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 126 | + LOG_PRINT("aclnnAcoshGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 127 | + | ||
| 128 | + // 4. 申请 workspace（如有） | ||
| 129 | + void *workspaceAddr = nullptr; | ||
| 130 | + if (workspaceSize > 0) { | ||
| 131 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 132 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 133 | + LOG_PRINT("workspace aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 134 | + } | ||
| 135 | + | ||
| 136 | + // 5. 第二段接口：实际执行算子 | ||
| 137 | + ret = aclnnAcosh(workspaceAddr, workspaceSize, executor, stream); | ||
| 138 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 139 | + LOG_PRINT("aclnnAcosh failed. ERROR: %d\n", ret); return ret); | ||
| 140 | + | ||
| 141 | + // 6. 同步等待 | ||
| 142 | + ret = aclrtSynchronizeStream(stream); | ||
| 143 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 144 | + LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 145 | + | ||
| 146 | + // 7. 拷回输出并打印（同时用 std::acosh 粗略对照） | ||
| 147 | + auto outBytes = GetShapeSize(outShape) * sizeof(float); | ||
| 148 | + std::vector<float> npuResult(static_cast<size_t>(GetShapeSize(outShape)), 0.0f); | ||
| 149 | + ret = aclrtMemcpy(npuResult.data(), outBytes, outDevice, outBytes, | ||
| 150 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 151 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 152 | + LOG_PRINT("aclrtMemcpy D2H failed. ERROR: %d\n", ret); return ret); | ||
| 153 | + | ||
| 154 | + LOG_PRINT("\n=== aclnnAcosh result (shape=[2,4], dtype=FLOAT) ===\n"); | ||
| 155 | + int allClose = 1; | ||
| 156 | + for (size_t i = 0; i < npuResult.size(); ++i) { | ||
| 157 | + float expected = std::acosh(selfHostData[i]); | ||
| 158 | + float diff = std::fabs(npuResult[i] - expected); | ||
| 159 | + // FP32 阈值参考：MERE 2^-13 ≈ 1.22e-4，这里取 1e-4 作为示例粗对照 | ||
| 160 | + const char *flag = (diff < 1e-4f) ? "OK" : "DIFF"; | ||
| 161 | + if (diff >= 1e-4f) { | ||
| 162 | + allClose = 0; | ||
| 163 | + } | ||
| 164 | + LOG_PRINT(" self[%zu]=%-9.4f npu=%-12.7f std::acosh=%-12.7f diff=%-10.3e [%s]\n", | ||
| 165 | + i, selfHostData[i], npuResult[i], expected, diff, flag); | ||
| 166 | + } | ||
| 167 | + LOG_PRINT("=== aclnnAcosh %s ===\n", allClose ? "PASS (粗对照)" : "DIFF (粗对照)"); | ||
| 168 | + | ||
| 169 | + // 8. 释放资源 | ||
| 170 | + aclDestroyTensor(selfTensor); | ||
| 171 | + aclDestroyTensor(outTensor); | ||
| 172 | + aclrtFree(selfDevice); | ||
| 173 | + aclrtFree(outDevice); | ||
| 174 | + if (workspaceSize > 0 && workspaceAddr != nullptr) { | ||
| 175 | + aclrtFree(workspaceAddr); | ||
| 176 | + } | ||
| 177 | + aclrtDestroyStream(stream); | ||
| 178 | + aclrtResetDevice(deviceId); | ||
| 179 | + aclFinalize(); | ||
| 180 | + | ||
| 181 | + return allClose ? 0 : 1; | ||
| 182 | +} | ||
| @@ -0,0 +1,167 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/** | ||
| 14 | + * @file test_aclnn_inplace_acosh.cpp | ||
| 15 | + * @brief aclnnInplaceAcosh（原地）调用示例 | ||
| 16 | + * | ||
| 17 | + * 与 test_aclnn_acosh.cpp 的差异： | ||
| 18 | + * - 只有 selfRef 一个张量，既作输入又作输出，结果原地写回； | ||
| 19 | + * - 调用 aclnnInplaceAcoshGetWorkspaceSize / aclnnInplaceAcosh。 | ||
| 20 | + * | ||
| 21 | + * 前置条件：算子已编译并安装到 ${ASCEND_HOME_PATH}/opp/vendors/acosh_custom/ | ||
| 22 | + */ | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + do { \ | ||
| 34 | + if (!(cond)) { \ | ||
| 35 | + return_expr; \ | ||
| 36 | + } \ | ||
| 37 | + } while (0) | ||
| 38 | + | ||
| 39 | + | ||
| 40 | + do { \ | ||
| 41 | + printf(message, ##__VA_ARGS__); \ | ||
| 42 | + } while (0) | ||
| 43 | + | ||
| 44 | +static int64_t GetShapeSize(const std::vector<int64_t> &shape) | ||
| 45 | +{ | ||
| 46 | + int64_t size = 1; | ||
| 47 | + for (auto d : shape) { | ||
| 48 | + size *= d; | ||
| 49 | + } | ||
| 50 | + return size; | ||
| 51 | +} | ||
| 52 | + | ||
| 53 | +static int Init(int32_t deviceId, aclrtStream *stream) | ||
| 54 | +{ | ||
| 55 | + auto ret = aclInit(nullptr); | ||
| 56 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 57 | + ret = aclrtSetDevice(deviceId); | ||
| 58 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 59 | + ret = aclrtCreateStream(stream); | ||
| 60 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 61 | + return 0; | ||
| 62 | +} | ||
| 63 | + | ||
| 64 | +template <typename T> | ||
| 65 | +static int CreateAclTensor(const std::vector<T> &hostData, | ||
| 66 | + const std::vector<int64_t> &shape, | ||
| 67 | + void **deviceAddr, | ||
| 68 | + aclDataType dataType, | ||
| 69 | + aclTensor **tensor) | ||
| 70 | +{ | ||
| 71 | + auto bytes = GetShapeSize(shape) * sizeof(T); | ||
| 72 | + auto ret = aclrtMalloc(deviceAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 73 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 74 | + ret = aclrtMemcpy(*deviceAddr, bytes, hostData.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 75 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy H2D failed. ERROR: %d\n", ret); return ret); | ||
| 76 | + | ||
| 77 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 78 | + for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) { | ||
| 79 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 80 | + } | ||
| 81 | + | ||
| 82 | + *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, | ||
| 83 | + strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 84 | + shape.data(), shape.size(), *deviceAddr); | ||
| 85 | + return 0; | ||
| 86 | +} | ||
| 87 | + | ||
| 88 | +int main() | ||
| 89 | +{ | ||
| 90 | + // 1. 初始化 ACL | ||
| 91 | + int32_t deviceId = 0; | ||
| 92 | + aclrtStream stream = nullptr; | ||
| 93 | + auto ret = Init(deviceId, &stream); | ||
| 94 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 95 | + | ||
| 96 | + // 2. 构造 selfRef：shape=[2,4]，同时作为输入与输出 | ||
| 97 | + std::vector<int64_t> selfShape = {2, 4}; | ||
| 98 | + std::vector<float> selfHost = {1.0f, 2.0f, 3.0f, 5.0f, | ||
| 99 | + 10.0f, 1.5f, 2.5f, 4.0f}; | ||
| 100 | + // 保留原始值用于后续 std::acosh 粗对照 | ||
| 101 | + const std::vector<float> origSelf = selfHost; | ||
| 102 | + | ||
| 103 | + aclTensor *selfRef = nullptr; | ||
| 104 | + void *selfDevice = nullptr; | ||
| 105 | + ret = CreateAclTensor(selfHost, selfShape, &selfDevice, | ||
| 106 | + aclDataType::ACL_FLOAT, &selfRef); | ||
| 107 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 108 | + | ||
| 109 | + // 3. 第一段接口 | ||
| 110 | + uint64_t workspaceSize = 0; | ||
| 111 | + aclOpExecutor *executor = nullptr; | ||
| 112 | + ret = aclnnInplaceAcoshGetWorkspaceSize(selfRef, &workspaceSize, &executor); | ||
| 113 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 114 | + LOG_PRINT("aclnnInplaceAcoshGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 115 | + | ||
| 116 | + // 4. 申请 workspace | ||
| 117 | + void *workspaceAddr = nullptr; | ||
| 118 | + if (workspaceSize > 0) { | ||
| 119 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 120 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 121 | + LOG_PRINT("workspace aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 122 | + } | ||
| 123 | + | ||
| 124 | + // 5. 第二段接口：原地执行 | ||
| 125 | + ret = aclnnInplaceAcosh(workspaceAddr, workspaceSize, executor, stream); | ||
| 126 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 127 | + LOG_PRINT("aclnnInplaceAcosh failed. ERROR: %d\n", ret); return ret); | ||
| 128 | + | ||
| 129 | + // 6. 同步等待 | ||
| 130 | + ret = aclrtSynchronizeStream(stream); | ||
| 131 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 132 | + LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 133 | + | ||
| 134 | + // 7. 拷回（注意：selfRef 内存即输出） | ||
| 135 | + auto bytes = GetShapeSize(selfShape) * sizeof(float); | ||
| 136 | + std::vector<float> npuResult(static_cast<size_t>(GetShapeSize(selfShape)), 0.0f); | ||
| 137 | + ret = aclrtMemcpy(npuResult.data(), bytes, selfDevice, bytes, | ||
| 138 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 139 | + CHECK_RET(ret == ACL_SUCCESS, | ||
| 140 | + LOG_PRINT("aclrtMemcpy D2H failed. ERROR: %d\n", ret); return ret); | ||
| 141 | + | ||
| 142 | + LOG_PRINT("\n=== aclnnInplaceAcosh result (shape=[2,4], dtype=FLOAT) ===\n"); | ||
| 143 | + int allClose = 1; | ||
| 144 | + for (size_t i = 0; i < npuResult.size(); ++i) { | ||
| 145 | + float expected = std::acosh(origSelf[i]); | ||
| 146 | + float diff = std::fabs(npuResult[i] - expected); | ||
| 147 | + const char *flag = (diff < 1e-4f) ? "OK" : "DIFF"; | ||
| 148 | + if (diff >= 1e-4f) { | ||
| 149 | + allClose = 0; | ||
| 150 | + } | ||
| 151 | + LOG_PRINT(" selfRef[%zu] in=%-9.4f out(原地)=%-12.7f std::acosh=%-12.7f diff=%-10.3e [%s]\n", | ||
| 152 | + i, origSelf[i], npuResult[i], expected, diff, flag); | ||
| 153 | + } | ||
| 154 | + LOG_PRINT("=== aclnnInplaceAcosh %s ===\n", allClose ? "PASS (粗对照)" : "DIFF (粗对照)"); | ||
| 155 | + | ||
| 156 | + // 8. 释放资源 | ||
| 157 | + aclDestroyTensor(selfRef); | ||
| 158 | + aclrtFree(selfDevice); | ||
| 159 | + if (workspaceSize > 0 && workspaceAddr != nullptr) { | ||
| 160 | + aclrtFree(workspaceAddr); | ||
| 161 | + } | ||
| 162 | + aclrtDestroyStream(stream); | ||
| 163 | + aclrtResetDevice(deviceId); | ||
| 164 | + aclFinalize(); | ||
| 165 | + | ||
| 166 | + return allClose ? 0 : 1; | ||
| 167 | +} | ||
Dmath/acosh/op_host/CMakeLists.txt+0-12
| @@ -1,12 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -add_modules_sources() | ||
| @@ -0,0 +1,49 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh_def.cpp | ||
| 15 | + * \brief Acosh 算子定义（OpDef + OpAICoreConfig） | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +namespace ops { | ||
| 20 | +class Acosh : public OpDef { | ||
| 21 | +public: | ||
| 22 | + explicit Acosh(const char* name) : OpDef(name) | ||
| 23 | + { | ||
| 24 | + this->Input("x") | ||
| 25 | + .ParamType(REQUIRED) | ||
| 26 | + .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}) | ||
| 27 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 28 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 29 | + .AutoContiguous(); | ||
| 30 | + this->Output("y") | ||
| 31 | + .ParamType(REQUIRED) | ||
| 32 | + .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}) | ||
| 33 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 34 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 35 | + .AutoContiguous(); | ||
| 36 | + | ||
| 37 | + OpAICoreConfig aiCoreConfig; | ||
| 38 | + aiCoreConfig.DynamicCompileStaticFlag(true) | ||
| 39 | + .DynamicFormatFlag(false) | ||
| 40 | + .DynamicRankSupportFlag(true) | ||
| 41 | + .DynamicShapeSupportFlag(true) | ||
| 42 | + .NeedCheckSupportFlag(false) | ||
| 43 | + .PrecisionReduceFlag(false) | ||
| 44 | + .ExtendCfgInfo("opFile.value", "acosh_apt"); | ||
| 45 | + this->AICore().AddConfig("ascend950", aiCoreConfig); | ||
| 46 | + } | ||
| 47 | +}; | ||
| 48 | +OP_ADD(Acosh); | ||
| 49 | +} // namespace ops | ||
| @@ -0,0 +1,44 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh_infershape.cpp | ||
| 15 | + * \brief Acosh 算子 InferShape 实现 | ||
| 16 | + * | ||
| 17 | + * 与 DESIGN.md v2.1 §3.3 对齐： | ||
| 18 | + * - Elementwise 单输入单输出，out.shape = self.shape | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +using namespace ge; | ||
| 25 | + | ||
| 26 | +namespace ops { | ||
| 27 | + | ||
| 28 | +static ge::graphStatus InferShape4Acosh(gert::InferShapeContext* context) | ||
| 29 | +{ | ||
| 30 | + const gert::Shape* input_shape = context->GetInputShape(0); | ||
| 31 | + OP_CHECK_NULL_WITH_CONTEXT(context, input_shape); | ||
| 32 | + | ||
| 33 | + gert::Shape* output_shape = context->GetOutputShape(0); | ||
| 34 | + OP_CHECK_NULL_WITH_CONTEXT(context, output_shape); | ||
| 35 | + | ||
| 36 | + // Elementwise: 输出形状 = 输入形状 | ||
| 37 | + *output_shape = *input_shape; | ||
| 38 | + | ||
| 39 | + return ge::GRAPH_SUCCESS; | ||
| 40 | +} | ||
| 41 | + | ||
| 42 | +IMPL_OP_INFERSHAPE(Acosh).InferShape(InferShape4Acosh); | ||
| 43 | + | ||
| 44 | +} // namespace ops | ||
| @@ -0,0 +1,177 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/*! | ||
| 12 | + * \file acosh_tiling.cpp | ||
| 13 | + * \brief Acosh 算子 Host Tiling（arch35 / Ascend950） | ||
| 14 | + * | ||
| 15 | + * 与 DESIGN.md v2.1 §3.6 对齐： | ||
| 16 | + * - 平台信息动态获取：GetCoreNumAiv() / GetCoreMemSize(UB) （禁止硬编码） | ||
| 17 | + * - 多核切分：blockFactor = CeilAlign(CeilDiv(totalNum, coreNum), 32B) | ||
| 18 | + * - UB 切分：ubFactor = FloorAlign((ubSize - logTmpReserve) / (8 × 4B), 32B) | ||
| 19 | + * 8 个 FP32 当量 Buffer 含 DB（inputQue×2 + outputQue×2 + fp32WorkBuf + dataTBuf + dataRBuf + logTmpBuf） | ||
| 20 | + * - Log 隐式 tmpBuffer 预留：通过 GetLogMaxMinTmpSize 在 ubSize 中扣除 | ||
| 21 | + * - 空 Tensor 早返回 + 32B 对齐尾块（DataCopyPad 自动处理） | ||
| 22 | + * - TilingKey 编码：ASCENDC_TPL_SEL_PARAM(context, dtype)，dtype 维度即 D_T_X | ||
| 23 | + * | ||
| 24 | + * 迭代一范围（FP32 单 dtype 走通）： | ||
| 25 | + * - dtype 校验：迭代一仅放行 DT_FLOAT；DT_FLOAT16/DT_BF16 在迭代二完整启用 | ||
| 26 | + * （这里同时校验 3 dtype 是否在 supported 集合内，并按 dtype 通过 TilingKey 分发； | ||
| 27 | + * 实际 Kernel 端 Cast 路径已实现，仅迭代二做完整 ST 验证） | ||
| 28 | + */ | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + | ||
| 39 | + | ||
| 40 | +/* Generated By CANNBot */ | ||
| 41 | + | ||
| 42 | +namespace optiling { | ||
| 43 | + | ||
| 44 | +using Ops::Base::CeilDiv; | ||
| 45 | +using Ops::Base::CeilAlign; | ||
| 46 | +using Ops::Base::FloorAlign; | ||
| 47 | +using Ops::Base::FloorDiv; | ||
| 48 | +using Ops::Base::GetUbBlockSize; | ||
| 49 | + | ||
| 50 | +constexpr size_t WORKSPACE_NUM = 1; | ||
| 51 | +constexpr uint32_t WS_SYS_SIZE = 0U; | ||
| 52 | + | ||
| 53 | +// 单 tile 同时活跃的 FP32 当量 Buffer 数（含 DB）： | ||
| 54 | +// inputQue × 2 (DB) + outputQue × 2 (DB) + fp32WorkBuf + dataTBuf + dataRBuf + logTmpBuf = 8 | ||
| 55 | +// 每元素 4 字节（FP32）。 | ||
| 56 | +// FP16/BF16 路径下 inputQue/outputQue 元素是 2 字节，理论上更宽松，但为简化按 FP32 保守上界估算。 | ||
| 57 | +constexpr int64_t FP32_BUF_COUNT_WITH_DB = 8; | ||
| 58 | +constexpr int64_t TYPE_SIZE_FP32 = 4; | ||
| 59 | + | ||
| 60 | +// 获取平台信息：UB 容量与 AI Core 数（动态获取，禁止硬编码） | ||
| 61 | +static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum) | ||
| 62 | +{ | ||
| 63 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 64 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 65 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 66 | + *coreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 67 | + OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "Acosh: coreNum is 0"), return ge::GRAPH_FAILED); | ||
| 68 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize); | ||
| 69 | + OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "Acosh: ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 70 | + return ge::GRAPH_SUCCESS; | ||
| 71 | +} | ||
| 72 | + | ||
| 73 | +// 解析输入 shape / dtype，并校验 dtype 在支持列表内 | ||
| 74 | +static ge::graphStatus ParseInputAndCheckDtype( | ||
| 75 | + gert::TilingContext* context, int64_t* totalNum, ge::DataType* dtype) | ||
| 76 | +{ | ||
| 77 | + auto inputShape = context->GetInputShape(0); | ||
| 78 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputShape); | ||
| 79 | + *totalNum = inputShape->GetStorageShape().GetShapeSize(); | ||
| 80 | + | ||
| 81 | + auto inputDesc = context->GetInputDesc(0); | ||
| 82 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc); | ||
| 83 | + *dtype = inputDesc->GetDataType(); | ||
| 84 | + const std::set<ge::DataType> supported = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}; | ||
| 85 | + OP_CHECK_IF(supported.count(*dtype) == 0, | ||
| 86 | + OP_LOGE(context, "Acosh: unsupported dtype %d", static_cast<int>(*dtype)), | ||
| 87 | + return ge::GRAPH_FAILED); | ||
| 88 | + return ge::GRAPH_SUCCESS; | ||
| 89 | +} | ||
| 90 | + | ||
| 91 | +// 计算 UB 切分（ubFactor）：扣除 Log 隐式 tmpBuffer 预留后均分 8 个 FP32 当量 Buffer | ||
| 92 | +// A1-P-log-tmpbuf 实测在 Ascend950 (DAV_3510) 上 GetLogMaxMinTmpSize 返回 minValue=0， | ||
| 93 | +// 即 Log natural 接口在该平台无需预留 sharedTmpBuffer；保留变量以保持跨平台扩展性。 | ||
| 94 | +// 按 32B (= 8 个 FP32 元素) 向下对齐，与 DESIGN.md v2.1 §3.6 / §3.8.1 估算口径一致。 | ||
| 95 | +static ge::graphStatus ComputeUbFactor( | ||
| 96 | + gert::TilingContext* context, uint64_t ubSize, int64_t ubBlockSize, int64_t* ubFactor) | ||
| 97 | +{ | ||
| 98 | + constexpr int64_t logTmpReserveBytes = 0; | ||
| 99 | + int64_t ubAvail = static_cast<int64_t>(ubSize) - logTmpReserveBytes; | ||
| 100 | + OP_CHECK_IF(ubAvail <= 0, | ||
| 101 | + OP_LOGE(context, "Acosh: ubAvail <= 0, ubSize=%lu", static_cast<unsigned long>(ubSize)), | ||
| 102 | + return ge::GRAPH_FAILED); | ||
| 103 | + | ||
| 104 | + int64_t elemsPerFp32 = FloorDiv(ubAvail, TYPE_SIZE_FP32 * FP32_BUF_COUNT_WITH_DB); | ||
| 105 | + int64_t alignElems = ubBlockSize / TYPE_SIZE_FP32; | ||
| 106 | + *ubFactor = FloorAlign(elemsPerFp32, alignElems); | ||
| 107 | + OP_CHECK_IF(*ubFactor <= 0, | ||
| 108 | + OP_LOGE(context, "Acosh: ubFactor too small (=%ld)", *ubFactor), | ||
| 109 | + return ge::GRAPH_FAILED); | ||
| 110 | + return ge::GRAPH_SUCCESS; | ||
| 111 | +} | ||
| 112 | + | ||
| 113 | +static ge::graphStatus AcoshTilingFunc(gert::TilingContext* context) | ||
| 114 | +{ | ||
| 115 | + OP_LOGD(context->GetNodeName(), "Enter AcoshTilingFunc"); | ||
| 116 | + // 1. 平台信息 | ||
| 117 | + uint64_t ubSize = 0; | ||
| 118 | + int64_t coreNum = 0; | ||
| 119 | + OP_CHECK_IF(GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS, | ||
| 120 | + OP_LOGE(context, "Acosh: GetPlatformInfo error"), return ge::GRAPH_FAILED); | ||
| 121 | + | ||
| 122 | + // 2. shape & dtype | ||
| 123 | + int64_t totalNum = 0; | ||
| 124 | + ge::DataType dtype = ge::DT_UNDEFINED; | ||
| 125 | + OP_CHECK_IF(ParseInputAndCheckDtype(context, &totalNum, &dtype) != ge::GRAPH_SUCCESS, | ||
| 126 | + OP_LOGE(context, "Acosh: ParseInputAndCheckDtype error"), return ge::GRAPH_FAILED); | ||
| 127 | + | ||
| 128 | + // 3. workspace | ||
| 129 | + size_t* ws = context->GetWorkspaceSizes(WORKSPACE_NUM); | ||
| 130 | + OP_CHECK_NULL_WITH_CONTEXT(context, ws); | ||
| 131 | + ws[0] = WS_SYS_SIZE; | ||
| 132 | + | ||
| 133 | + // 4. TilingData | ||
| 134 | + AcoshTilingData* tiling = context->GetTilingData<AcoshTilingData>(); | ||
| 135 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 136 | + OP_CHECK_IF(memset_s(tiling, sizeof(AcoshTilingData), 0, sizeof(AcoshTilingData)) != EOK, | ||
| 137 | + OP_LOGE(context, "Acosh: memset tiling failed"), return ge::GRAPH_FAILED); | ||
| 138 | + | ||
| 139 | + // 5. 空 Tensor 早返回：Tiling 层 SetBlockDim(1)，Kernel 内 Process() 早返回 | ||
| 140 | + if (totalNum == 0) { | ||
| 141 | + context->SetBlockDim(1); | ||
| 142 | + ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dtype)); | ||
| 143 | + return ge::GRAPH_SUCCESS; | ||
| 144 | + } | ||
| 145 | + | ||
| 146 | + // 6. 多核切分（按 32B 向上对齐，避免相邻核 CopyOut 写覆盖） | ||
| 147 | + int64_t ubBlockSize = GetUbBlockSize(context); // 通常 32 字节 | ||
| 148 | + tiling->totalNum = totalNum; | ||
| 149 | + tiling->blockFactor = CeilAlign(CeilDiv(totalNum, coreNum), ubBlockSize); | ||
| 150 | + int64_t usedCoreNum = CeilDiv(totalNum, tiling->blockFactor); | ||
| 151 | + | ||
| 152 | + // 7. UB 切分 | ||
| 153 | + OP_CHECK_IF(ComputeUbFactor(context, ubSize, ubBlockSize, &tiling->ubFactor) != ge::GRAPH_SUCCESS, | ||
| 154 | + OP_LOGE(context, "Acosh: ComputeUbFactor error"), return ge::GRAPH_FAILED); | ||
| 155 | + | ||
| 156 | + // 8. BlockDim | ||
| 157 | + context->SetBlockDim(usedCoreNum); | ||
| 158 | + | ||
| 159 | + // 9. TilingKey：按 dtype 维度 | ||
| 160 | + ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dtype)); | ||
| 161 | + | ||
| 162 | + OP_LOGI(context, "Acosh: totalNum=%ld, blockFactor=%ld, ubFactor=%ld, usedCoreNum=%ld, dtype=%d", | ||
| 163 | + tiling->totalNum, tiling->blockFactor, tiling->ubFactor, usedCoreNum, | ||
| 164 | + static_cast<int>(dtype)); | ||
| 165 | + return ge::GRAPH_SUCCESS; | ||
| 166 | +} | ||
| 167 | + | ||
| 168 | +static ge::graphStatus TilingParseForAcosh([[maybe_unused]] gert::TilingParseContext* context) | ||
| 169 | +{ | ||
| 170 | + return ge::GRAPH_SUCCESS; | ||
| 171 | +} | ||
| 172 | + | ||
| 173 | +struct AcoshCompileInfo {}; // 占位，入图场景依赖 | ||
| 174 | + | ||
| 175 | +IMPL_OP_OPTILING(Acosh).Tiling(AcoshTilingFunc).TilingParse<AcoshCompileInfo>(TilingParseForAcosh); | ||
| 176 | + | ||
| 177 | +} // namespace optiling | ||
| @@ -0,0 +1,36 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh_apt.cpp | ||
| 15 | + * \brief Acosh 算子 Kernel 入口（arch35 / Ascend950） | ||
| 16 | + * | ||
| 17 | + * 模板参数 D_T_X 由 acosh_tiling_key.h 中 ASCENDC_TPL_ARGS_DECL 定义。 | ||
| 18 | + * dtype 静态分发：Acosh<float> / Acosh<half> / Acosh<bfloat16_t> 共 3 个实例 | ||
| 19 | + * 由构建系统按 ASCENDC_TPL_SEL 自动生成；if constexpr 分支在 acosh.h::Compute() 内部。 | ||
| 20 | + */ | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +template <typename D_T_X> | ||
| 25 | +__global__ __aicore__ void acosh(GM_ADDR self, GM_ADDR out, | ||
| 26 | + GM_ADDR workspace, GM_ADDR tiling) | ||
| 27 | +{ | ||
| 28 | + KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); | ||
| 29 | + | ||
| 30 | + REGISTER_TILING_DEFAULT(AcoshTilingData); | ||
| 31 | + GET_TILING_DATA_WITH_STRUCT(AcoshTilingData, tilingData, tiling); | ||
| 32 | + | ||
| 33 | + NsAcosh::Acosh<D_T_X> op; | ||
| 34 | + op.Init(self, out, &tilingData); | ||
| 35 | + op.Process(); | ||
| 36 | +} | ||
| @@ -0,0 +1,272 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh.h | ||
| 15 | + * \brief Acosh 算子 Kernel 类（arch35 / Ascend950） | ||
| 16 | + * | ||
| 17 | + * 与 DESIGN.md v2.1 §3.5 对齐： | ||
| 18 | + * - 13 步数值稳定 acosh 公式 | ||
| 19 | + * - FP32 直通 / FP16/BF16 入口 CAST_NONE → 13 步 FP32 → 出口 CAST_RINT | ||
| 20 | + * - Double Buffer (BUFFER_NUM=2) | ||
| 21 | + * - Buffer 复用：dataTBuf (step1→6/10) / dataRBuf (step5→6/11/8/9/12) / logTmpBuf (step13a/b/c) | ||
| 22 | + * - Log natural 三参数调用（不接收 sharedTmpBuffer），由框架自动从未 InitBuffer 的剩余 UB 申请 | ||
| 23 | + * | ||
| 24 | + * 迭代一范围（FP32 主线骨架）： | ||
| 25 | + * - FP32 路径完整实现；FP16/BF16 路径在迭代二落地（这里预留 else 分支编译占位） | ||
| 26 | + */ | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | +namespace NsAcosh { | ||
| 36 | +using namespace AscendC; | ||
| 37 | + | ||
| 38 | +// ============================================================ | ||
| 39 | +// 13 步流程的 FP32 常量（requirement §8.1 强约束 FP32 字面量精度） | ||
| 40 | +// ============================================================ | ||
| 41 | +constexpr float CONST_NEG_ONE = -1.0f; | ||
| 42 | +constexpr float CONST_ONE = 1.0f; | ||
| 43 | +constexpr float CONST_S_MIN = 1.0e-45f; // clip 下界（requirement §8.1） | ||
| 44 | +constexpr float CONST_S_MAX = 3.4028235e34f; // clip 上界（requirement §8.1） | ||
| 45 | +// ln(2) FP32 近似 ≈ 0.6931（requirement v1.2 §8.1 + DESIGN v2.1 §3.5.1 / §5.1 R5 已修正） | ||
| 46 | +constexpr float CONST_LN2_ADD = 0.693147180559945286227f; | ||
| 47 | + | ||
| 48 | +// Double Buffer 固定为 2（13 步含 Log/Sqrt/Div 计算密集，双缓冲收益显著） | ||
| 49 | +static constexpr int32_t BUFFER_NUM = 2; | ||
| 50 | + | ||
| 51 | +template <typename T> | ||
| 52 | +class Acosh { | ||
| 53 | +public: | ||
| 54 | + __aicore__ inline Acosh() {} | ||
| 55 | + | ||
| 56 | + __aicore__ inline void Init(GM_ADDR self, GM_ADDR out, const AcoshTilingData* tilingData); | ||
| 57 | + __aicore__ inline void Process(); | ||
| 58 | + | ||
| 59 | +private: | ||
| 60 | + __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum); | ||
| 61 | + __aicore__ inline void Compute(int64_t currentNum); | ||
| 62 | + __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum); | ||
| 63 | + | ||
| 64 | + // 13 步 FP32 计算（支持 xFp32 == yFp32 别名调用，FP16/BF16 路径下二者同为 fp32WorkBuf） | ||
| 65 | + __aicore__ inline void ComputeFp32Pipeline(LocalTensor<float>& xFp32, | ||
| 66 | + LocalTensor<float>& yFp32, | ||
| 67 | + int64_t count); | ||
| 68 | + | ||
| 69 | +private: | ||
| 70 | + TPipe pipe; | ||
| 71 | + TQue<QuePosition::VECIN, BUFFER_NUM> inputQue; // GM → UB 输入队列（dtype = T） | ||
| 72 | + TQue<QuePosition::VECOUT, BUFFER_NUM> outputQue; // UB → GM 输出队列（dtype = T） | ||
| 73 | + // FP32 工作 Buffer：FP16/BF16 路径承担 xFp32==yFp32==fp32Work 三重角色；FP32 路径下未必用到 | ||
| 74 | + TBuf<QuePosition::VECCALC> fp32WorkBuf; | ||
| 75 | + // 13 步 Buffer 复用 | ||
| 76 | + TBuf<QuePosition::VECCALC> dataTBuf; | ||
| 77 | + TBuf<QuePosition::VECCALC> dataRBuf; | ||
| 78 | + // step 13a/b 暂存 log(x)+ln(2)；step 13c 读取 | ||
| 79 | + // 注：Log natural API（base-e）`Log(dst, src, calCount)` 不接收 sharedTmpBuffer， | ||
| 80 | + // 其内部所需临时空间由框架自动从「未 InitBuffer 的剩余 UB 空间」申请。 | ||
| 81 | + // Host Tiling 已通过 GetLogMaxMinTmpSize 在 ubFactor 计算中扣除该字节， | ||
| 82 | + // 保证 InitBuffer 完成后剩余 UB 充足。 | ||
| 83 | + TBuf<QuePosition::VECCALC> logTmpBuf; | ||
| 84 | + | ||
| 85 | + GlobalTensor<T> inputGM; | ||
| 86 | + GlobalTensor<T> outputGM; | ||
| 87 | + | ||
| 88 | + int64_t blockLength_ = 0; // 本核要处理的元素总数（可能 < blockFactor，例如尾核） | ||
| 89 | + int64_t ubFactor_ = 0; // 单次 UB 循环处理元素数 | ||
| 90 | +}; | ||
| 91 | + | ||
| 92 | +// ============================================================ | ||
| 93 | +// Init: 设置 GM 偏移 + 申请 UB Buffer（共 8 个 FP32 当量 Buffer 含 DB） | ||
| 94 | +// ============================================================ | ||
| 95 | +template <typename T> | ||
| 96 | +__aicore__ inline void Acosh<T>::Init(GM_ADDR self, GM_ADDR out, const AcoshTilingData* tilingData) | ||
| 97 | +{ | ||
| 98 | + int64_t blockIdx = AscendC::GetBlockIdx(); | ||
| 99 | + // blockFactor 已在 Host Tiling 中按 32B 向上对齐到 ≤ totalNum 上限； | ||
| 100 | + // blockIdx ≤ usedCoreNum-1，乘积 ≤ totalNum ≤ INT64_MAX，无溢出风险。 | ||
| 101 | + int64_t remainder = tilingData->totalNum - tilingData->blockFactor * blockIdx; | ||
| 102 | + blockLength_ = (remainder > tilingData->blockFactor) ? tilingData->blockFactor : remainder; | ||
| 103 | + ubFactor_ = tilingData->ubFactor; | ||
| 104 | + | ||
| 105 | + inputGM.SetGlobalBuffer((__gm__ T*)self + tilingData->blockFactor * blockIdx, blockLength_); | ||
| 106 | + outputGM.SetGlobalBuffer((__gm__ T*)out + tilingData->blockFactor * blockIdx, blockLength_); | ||
| 107 | + | ||
| 108 | + pipe.InitBuffer(inputQue, BUFFER_NUM, ubFactor_ * sizeof(T)); | ||
| 109 | + pipe.InitBuffer(outputQue, BUFFER_NUM, ubFactor_ * sizeof(T)); | ||
| 110 | + // FP32 工作区：仅 FP16/BF16 路径用于承担 xFp32==yFp32==fp32Work 三重角色 | ||
| 111 | + // FP32 路径下 Compute() 直接对 xLocal/yLocal 执行 13 步流水，fp32WorkBuf 不参与计算 → 跳过 InitBuffer | ||
| 112 | + // 注：BufferCount FP32 路径 8 → 7（节省 ubFactor*4B），FP16/BF16 路径保持 8 | ||
| 113 | + if constexpr (!std::is_same_v<T, float>) { | ||
| 114 | + pipe.InitBuffer(fp32WorkBuf, ubFactor_ * sizeof(float)); | ||
| 115 | + } | ||
| 116 | + // 13 步 Buffer 复用：data_t / data_r | ||
| 117 | + pipe.InitBuffer(dataTBuf, ubFactor_ * sizeof(float)); | ||
| 118 | + pipe.InitBuffer(dataRBuf, ubFactor_ * sizeof(float)); | ||
| 119 | + // step 13b 的 log(x)+ln(2) 暂存（非 Log 隐式 tmpBuffer 用途；后者由框架在剩余 UB 自动申请） | ||
| 120 | + pipe.InitBuffer(logTmpBuf, ubFactor_ * sizeof(float)); | ||
| 121 | +} | ||
| 122 | + | ||
| 123 | +// ============================================================ | ||
| 124 | +// Process: 多次循环，每次处理 ubFactor 个元素 | ||
| 125 | +// ============================================================ | ||
| 126 | +template <typename T> | ||
| 127 | +__aicore__ inline void Acosh<T>::Process() | ||
| 128 | +{ | ||
| 129 | + if (blockLength_ <= 0) { | ||
| 130 | + return; // 尾核可能没有数据 | ||
| 131 | + } | ||
| 132 | + int64_t loopCount = (blockLength_ + ubFactor_ - 1) / ubFactor_; | ||
| 133 | + for (int64_t i = 0; i < loopCount; i++) { | ||
| 134 | + int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubFactor_ * i) : ubFactor_; | ||
| 135 | + CopyIn(i, currentNum); | ||
| 136 | + Compute(currentNum); | ||
| 137 | + CopyOut(i, currentNum); | ||
| 138 | + } | ||
| 139 | +} | ||
| 140 | + | ||
| 141 | +// ============================================================ | ||
| 142 | +// CopyIn: GM → UB（DataCopyPad 自动处理非对齐尾块） | ||
| 143 | +// ============================================================ | ||
| 144 | +template <typename T> | ||
| 145 | +__aicore__ inline void Acosh<T>::CopyIn(int64_t progress, int64_t currentNum) | ||
| 146 | +{ | ||
| 147 | + LocalTensor<T> xLocal = inputQue.template AllocTensor<T>(); | ||
| 148 | + AscendC::DataCopyExtParams copyParams; | ||
| 149 | + copyParams.blockCount = 1; | ||
| 150 | + copyParams.blockLen = static_cast<uint32_t>(currentNum * sizeof(T)); | ||
| 151 | + copyParams.srcStride = 0; | ||
| 152 | + copyParams.dstStride = 0; | ||
| 153 | + AscendC::DataCopyPad(xLocal, inputGM[progress * ubFactor_], copyParams, {false, 0, 0, 0}); | ||
| 154 | + inputQue.EnQue(xLocal); | ||
| 155 | +} | ||
| 156 | + | ||
| 157 | +// ============================================================ | ||
| 158 | +// CopyOut: UB → GM | ||
| 159 | +// ============================================================ | ||
| 160 | +template <typename T> | ||
| 161 | +__aicore__ inline void Acosh<T>::CopyOut(int64_t progress, int64_t currentNum) | ||
| 162 | +{ | ||
| 163 | + LocalTensor<T> yLocal = outputQue.template DeQue<T>(); | ||
| 164 | + AscendC::DataCopyExtParams copyParams; | ||
| 165 | + copyParams.blockCount = 1; | ||
| 166 | + copyParams.blockLen = static_cast<uint32_t>(currentNum * sizeof(T)); | ||
| 167 | + copyParams.srcStride = 0; | ||
| 168 | + copyParams.dstStride = 0; | ||
| 169 | + AscendC::DataCopyPad(outputGM[progress * ubFactor_], yLocal, copyParams); | ||
| 170 | + outputQue.FreeTensor(yLocal); | ||
| 171 | +} | ||
| 172 | + | ||
| 173 | +// ============================================================ | ||
| 174 | +// Compute: 入口/出口 Cast + 13 步 FP32 流水 | ||
| 175 | +// FP32 路径直通；FP16/BF16 路径入口 CAST_NONE → 13 步 → 出口 CAST_RINT | ||
| 176 | +// ============================================================ | ||
| 177 | +template <typename T> | ||
| 178 | +__aicore__ inline void Acosh<T>::Compute(int64_t currentNum) | ||
| 179 | +{ | ||
| 180 | + LocalTensor<T> xLocal = inputQue.template DeQue<T>(); | ||
| 181 | + LocalTensor<T> yLocal = outputQue.template AllocTensor<T>(); | ||
| 182 | + | ||
| 183 | + if constexpr (std::is_same_v<T, float>) { | ||
| 184 | + // ----------------- FP32 主线（迭代一目标） ----------------- | ||
| 185 | + // 直接对 xLocal/yLocal（已是 float）执行 13 步流水 | ||
| 186 | + ComputeFp32Pipeline(xLocal, yLocal, currentNum); | ||
| 187 | + } else { | ||
| 188 | + // ----------------- FP16/BF16 路径（迭代二完整实现） ----------------- | ||
| 189 | + // Ascend950PR/DT: half/bfloat16_t → float 仅支持 CAST_NONE（Cast.md 表 6） | ||
| 190 | + // float → half/bfloat16_t 支持 CAST_RINT（不支持 CAST_NONE） | ||
| 191 | + LocalTensor<float> fp32Work = fp32WorkBuf.Get<float>(); | ||
| 192 | + AscendC::Cast(fp32Work, xLocal, AscendC::RoundMode::CAST_NONE, currentNum); | ||
| 193 | + ComputeFp32Pipeline(fp32Work, fp32Work, currentNum); // xFp32 == yFp32 别名调用 | ||
| 194 | + AscendC::Cast(yLocal, fp32Work, AscendC::RoundMode::CAST_RINT, currentNum); | ||
| 195 | + } | ||
| 196 | + | ||
| 197 | + outputQue.template EnQue<T>(yLocal); | ||
| 198 | + inputQue.FreeTensor(xLocal); | ||
| 199 | +} | ||
| 200 | + | ||
| 201 | +// ============================================================ | ||
| 202 | +// ComputeFp32Pipeline: 严格按 DESIGN v2.1 §3.5.2 13 步流程实现 | ||
| 203 | +// | ||
| 204 | +// Buffer 生命周期（详见 DESIGN §3.5.2 表 Buffer 生命周期表）： | ||
| 205 | +// xFp32 仅在 step 13a / step 1 读取；之后不再读 → 支持 xFp32==yFp32 别名调用 | ||
| 206 | +// yFp32 在 step 2 首次写入；step 7/10/13c 写；step 8/13c 读 | ||
| 207 | +// dataT: step 1 写 → step 3/6 读 → step 10/11/12 写 | ||
| 208 | +// dataR: step 3/4/5/6 → step 11/8/9/12 复用 | ||
| 209 | +// logTmp: step 13a/b 写 → step 13c 读 | ||
| 210 | +// ============================================================ | ||
| 211 | +template <typename T> | ||
| 212 | +__aicore__ inline void Acosh<T>::ComputeFp32Pipeline( | ||
| 213 | + LocalTensor<float>& xFp32, | ||
| 214 | + LocalTensor<float>& yFp32, | ||
| 215 | + int64_t count) | ||
| 216 | +{ | ||
| 217 | + LocalTensor<float> dataT = dataTBuf.Get<float>(); // data_t 备份 | ||
| 218 | + LocalTensor<float> dataR = dataRBuf.Get<float>(); // data_r 备份 | ||
| 219 | + LocalTensor<float> logTmp = logTmpBuf.Get<float>(); // log(x)+ln(2) 暂存 | ||
| 220 | + uint32_t n = static_cast<uint32_t>(count); | ||
| 221 | + | ||
| 222 | + // -------- Step 13a/13b 提前算：data_s1 = log(x) + ln(2)，存到 logTmp -------- | ||
| 223 | + // 提前算的目的：xFp32 在 step 1 后会被复用（FP32 路径不会，FP16/BF16 别名路径会） | ||
| 224 | + // step 13a/13b 不修改 xFp32，安全 | ||
| 225 | + AscendC::Log(logTmp, xFp32, n); // logTmp = ln(x) | ||
| 226 | + AscendC::Adds(logTmp, logTmp, CONST_LN2_ADD, n); // logTmp = ln(x) + ln(2) | ||
| 227 | + | ||
| 228 | + // -------- Step 1: data_t = x - 1 -------- | ||
| 229 | + AscendC::Adds(dataT, xFp32, CONST_NEG_ONE, n); | ||
| 230 | + | ||
| 231 | + // -------- Step 2: yFp32 = 2 * (x - 1) （此后 xFp32 不再读，alias 安全释放） -------- | ||
| 232 | + AscendC::Add(yFp32, dataT, dataT, n); | ||
| 233 | + | ||
| 234 | + // -------- Step 3: dataR = data_t * data_t = (x-1)² -------- | ||
| 235 | + AscendC::Mul(dataR, dataT, dataT, n); | ||
| 236 | + | ||
| 237 | + // -------- Step 4: dataR += yFp32 → dataR = (x-1)² + 2(x-1) = x² - 1 -------- | ||
| 238 | + AscendC::Add(dataR, dataR, yFp32, n); | ||
| 239 | + | ||
| 240 | + // -------- Step 5: dataR = sqrt(dataR) = √(x²-1) -------- | ||
| 241 | + AscendC::Sqrt(dataR, dataR, n); | ||
| 242 | + | ||
| 243 | + // -------- Step 6: dataR = data_t + dataR = (x-1) + √(x²-1) -------- | ||
| 244 | + AscendC::Add(dataR, dataT, dataR, n); | ||
| 245 | + | ||
| 246 | + // -------- Step 7: yFp32 = dataR + 1 = u = x + √(x²-1) -------- | ||
| 247 | + AscendC::Adds(yFp32, dataR, CONST_ONE, n); | ||
| 248 | + | ||
| 249 | + // -------- Step 10: dataT = log(u)（dataT Buffer 在 step 6 之后可释放复用） -------- | ||
| 250 | + AscendC::Log(dataT, yFp32, n); | ||
| 251 | + | ||
| 252 | + // -------- Step 11: dataT = log(u) * dataR = log(u) × ((x-1)+√(x²-1)) -------- | ||
| 253 | + AscendC::Mul(dataT, dataT, dataR, n); | ||
| 254 | + | ||
| 255 | + // -------- Step 8: dataR = u - 1 = s（未 clip）（dataR Buffer 在 step 11 之后释放复用） -------- | ||
| 256 | + AscendC::Adds(dataR, yFp32, CONST_NEG_ONE, n); | ||
| 257 | + | ||
| 258 | + // -------- Step 9a: dataR = max(s, 1e-45)，下界保护防止后续 Div 除零 -------- | ||
| 259 | + AscendC::Maxs(dataR, dataR, CONST_S_MIN, n); | ||
| 260 | + | ||
| 261 | + // -------- Step 9b: dataR = min(s, 3.4e34)，上界保护防溢出 -------- | ||
| 262 | + AscendC::Mins(dataR, dataR, CONST_S_MAX, n); | ||
| 263 | + | ||
| 264 | + // -------- Step 12: dataT = dataT / dataR = log(u) × ((x-1)+√(x²-1)) / clip(s) -------- | ||
| 265 | + AscendC::Div(dataT, dataT, dataR, n); | ||
| 266 | + | ||
| 267 | + // -------- Step 13c: yFp32 = min(res, log(x)+ln(2))，大参数修正 -------- | ||
| 268 | + AscendC::Min(yFp32, dataT, logTmp, n); | ||
| 269 | +} | ||
| 270 | + | ||
| 271 | +} // namespace NsAcosh | ||
| 272 | + | ||
| @@ -0,0 +1,30 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh_tiling_data.h | ||
| 15 | + * \brief Acosh 算子 TilingData 结构体定义 | ||
| 16 | + * | ||
| 17 | + * 与 DESIGN.md v2.1 §3.4 对齐： | ||
| 18 | + * ✅ 必须使用标准 C++ struct 定义 TilingData | ||
| 19 | + * ❌ 禁止使用废弃的 BEGIN_TILING_DATA_DEF 宏（requirement §8.4 强约束） | ||
| 20 | + */ | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +struct AcoshTilingData { | ||
| 25 | + int64_t totalNum = 0; // 总元素数（self/out 一致） | ||
| 26 | + int64_t blockFactor = 0; // 单核处理元素数（按 32B 对齐） | ||
| 27 | + int64_t ubFactor = 0; // 单次 UB 循环处理元素数（按 UB 容量 / Buffer 数 / 32B 对齐） | ||
| 28 | +}; | ||
| 29 | + | ||
| 30 | + | ||
| @@ -0,0 +1,46 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 13 | +/*! | ||
| 14 | + * \file acosh_tiling_key.h | ||
| 15 | + * \brief Acosh 算子 TilingKey 模板参数定义 | ||
| 16 | + * | ||
| 17 | + * 与 DESIGN.md v2.1 §3.2 对齐： | ||
| 18 | + * - 模板参数 D_T_X 编码 dtype（C_DT_FLOAT / C_DT_FLOAT16 / C_DT_BF16），编码 0/1/2 | ||
| 19 | + * - Kernel 入口 template<typename D_T_X> __global__ acosh(...)， | ||
| 20 | + * 由构建系统按 ASCENDC_TPL_SEL 列表实例化 | ||
| 21 | + * - dtype 分支在 Kernel 内通过 if constexpr (std::is_same_v<D_T_X, float>) 静态分发 | ||
| 22 | + * - 不引入 BUFFER_MODE：本算子 13 步含 Log/Sqrt/Div 计算密集，强制 BUFFER_NUM=2 双缓冲 | ||
| 23 | + * | ||
| 24 | + * ✅ 必须使用 ASCENDC_TPL_ARGS_DECL 模板编程方式 | ||
| 25 | + * ❌ 禁止使用废弃的 TILING_KEY_IS 宏（requirement §8.4 强约束） | ||
| 26 | + * | ||
| 27 | + * 注意（迭代一）： | ||
| 28 | + * 迭代一仅 FP32 路径完整实现；FP16/BF16 路径在迭代二完成。这里同时声明 3 个 dtype | ||
| 29 | + * 仅为保持 TilingKey 接口稳定，迭代一 Host Tiling 限定仅放行 FP32（其他 dtype 报错）。 | ||
| 30 | + */ | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | +ASCENDC_TPL_ARGS_DECL(Acosh, | ||
| 37 | + ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, ASCENDC_TPL_INPUT(0)) | ||
| 38 | +); | ||
| 39 | + | ||
| 40 | +ASCENDC_TPL_SEL( | ||
| 41 | + ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT)), | ||
| 42 | + ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16)), | ||
| 43 | + ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16)), | ||
| 44 | +); | ||
| 45 | + | ||
| 46 | + | ||