| @@ -66,6 +66,16 @@ | |||
| 66 | <td>AI CPU</td> | 66 | <td>AI CPU</td> |
| 67 | <td>从输入图像中提取多个裁剪区域,并将它们统一调整为指定大小,支持双线性插值和最近邻插值。</td> | 67 | <td>从输入图像中提取多个裁剪区域,并将它们统一调整为指定大小,支持双线性插值和最近邻插值。</td> |
| 68 | </tr> | 68 | </tr> |
| 69 | + <tr> | ||
| 70 | + <td>image</td> | ||
| 71 | + <td><a href="../../image/extract_glimpse_v2/README.md">extract_glimpse_v2</a></td> | ||
| 72 | + <td>✓</td> | ||
| 73 | + <td>✓</td> | ||
| 74 | + <td>✗</td> | ||
| 75 | + <td>✓</td> | ||
| 76 | + <td>AI Core</td> | ||
| 77 | + <td>从批量输入图像中提取指定位置和大小的子图像(glimpse)。</td> | ||
| 78 | + </tr> | ||
| 69 | <tr> | 79 | <tr> |
| 70 | <td>image</td> | 80 | <td>image</td> |
| 71 | <td><a href="../../image/grid_sample/README.md">grid_sample</a></td> | 81 | <td><a href="../../image/grid_sample/README.md">grid_sample</a></td> |
| @@ -0,0 +1,27 @@ | |||
| 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 | +file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 14 | + | ||
| 15 | +if(NOT ENABLE_TEST AND NOT BENCHMARK) | ||
| 16 | + list(REMOVE_ITEM CURRENT_DIRS tests) | ||
| 17 | +endif() | ||
| 18 | + | ||
| 19 | +foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 20 | + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 21 | + add_subdirectory(${SUB_DIR}) | ||
| 22 | + endif() | ||
| 23 | +endforeach() | ||
| 24 | + | ||
| 25 | +set(SUPPORT_COMPUTE_UNIT "ascend950") | ||
| 26 | +set(SUPPORT_TILING_DIR "arch35") | ||
| 27 | +add_all_modules_sources(OPTYPE extract_glimpse_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE) | ||
| @@ -0,0 +1,125 @@ | |||
| 1 | +# ExtractGlimpseV2 | ||
C | |||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +| 产品 | 是否支持 | | ||
| 6 | +| :----------------------------------------------------------- | :------: | | ||
| 7 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | | ||
| 8 | +| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | × | | ||
| 9 | +| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | × | | ||
| 10 | +| <term>Atlas 200I/500 A2 推理产品</term> | × | | ||
| 11 | +| <term>Atlas 推理系列产品</term> | × | | ||
| 12 | +| <term>Atlas 训练系列产品</term> | × | | ||
| 13 | + | ||
| 14 | +## 功能说明 | ||
| 15 | + | ||
| 16 | +- 算子功能:从批量输入图像中提取指定位置和大小的子图像(glimpse)。根据给定的偏移坐标和裁剪尺寸,从每张输入图像中裁剪出一个子区域,输出为一批裁剪后的图像。 | ||
| 17 | + | ||
| 18 | +- 计算公式: | ||
| 19 | + | ||
| 20 | +$$ | ||
| 21 | +\text{对于每个 batch item } i: | ||
| 22 | +$$ | ||
| 23 | +$$ | ||
| 24 | +\text{1. 坐标变换: } offset\_y, offset\_x \text{ (根据 normalized/centered 属性)} | ||
| 25 | +$$ | ||
| 26 | +$$ | ||
| 27 | +\text{2. 裁剪区域: } [start\_y, end\_y) \times [start\_x, end\_x) \text{ (clamp 到图像边界)} | ||
| 28 | +$$ | ||
| 29 | +$$ | ||
| 30 | +\text{3. 数据拷贝: } glimpse[i][base\_y:base\_y+copy\_h][base\_x:base\_x+copy\_w][:] = input[i][start\_y:end\_y][start\_x:end\_x][:] | ||
| 31 | +$$ | ||
| 32 | +$$ | ||
| 33 | +\text{4. 越界填充: } glimpse \text{ 中未拷贝区域填零 (noise="zero")} | ||
| 34 | +$$ | ||
| 35 | + | ||
| 36 | +## 参数说明 | ||
| 37 | + | ||
| 38 | +<table style="undefined;table-layout: fixed; width: 980px"><colgroup> | ||
| 39 | + <col style="width: 100px"> | ||
| 40 | + <col style="width: 150px"> | ||
| 41 | + <col style="width: 280px"> | ||
| 42 | + <col style="width: 330px"> | ||
| 43 | + <col style="width: 120px"> | ||
| 44 | + </colgroup> | ||
| 45 | + <thead> | ||
| 46 | + <tr> | ||
| 47 | + <th>参数名</th> | ||
| 48 | + <th>输入/输出/属性</th> | ||
| 49 | + <th>描述</th> | ||
| 50 | + <th>数据类型</th> | ||
| 51 | + <th>数据格式</th> | ||
| 52 | + </tr></thead> | ||
| 53 | + <tbody> | ||
| 54 | + <tr> | ||
| 55 | + <td>input</td> | ||
| 56 | + <td>输入</td> | ||
| 57 | + <td>输入图像批次,4D tensor,shape为[batch, height, width, channels]。</td> | ||
| 58 | + <td>FLOAT</td> | ||
| 59 | + <td>ND</td> | ||
| 60 | + </tr> | ||
| 61 | + <tr> | ||
| 62 | + <td>size</td> | ||
| 63 | + <td>输入</td> | ||
| 64 | + <td>裁剪尺寸[glimpse_h, glimpse_w],1D const tensor,shape为[2]。</td> | ||
| 65 | + <td>INT32</td> | ||
| 66 | + <td>ND</td> | ||
| 67 | + </tr> | ||
| 68 | + <tr> | ||
| 69 | + <td>offsets</td> | ||
| 70 | + <td>输入</td> | ||
| 71 | + <td>每个batch item的偏移坐标(y, x),2D tensor,shape为[batch, 2]。</td> | ||
| 72 | + <td>FLOAT</td> | ||
| 73 | + <td>ND</td> | ||
| 74 | + </tr> | ||
| 75 | + <tr> | ||
| 76 | + <td>glimpse</td> | ||
| 77 | + <td>输出</td> | ||
| 78 | + <td>裁剪后的图像批次,4D tensor,shape为[batch, size_h, size_w, channels]。</td> | ||
| 79 | + <td>FLOAT</td> | ||
| 80 | + <td>ND</td> | ||
| 81 | + </tr> | ||
| 82 | + <tr> | ||
| 83 | + <td>centered</td> | ||
| 84 | + <td>属性</td> | ||
| 85 | + <td>偏移坐标是否相对于图像居中。默认值:true。</td> | ||
| 86 | + <td>Bool</td> | ||
| 87 | + <td>-</td> | ||
| 88 | + </tr> | ||
| 89 | + <tr> | ||
| 90 | + <td>normalized</td> | ||
| 91 | + <td>属性</td> | ||
| 92 | + <td>偏移坐标是否归一化到[-1, 1]。默认值:true。</td> | ||
| 93 | + <td>Bool</td> | ||
| 94 | + <td>-</td> | ||
| 95 | + </tr> | ||
| 96 | + <tr> | ||
| 97 | + <td>uniform_noise</td> | ||
| 98 | + <td>属性</td> | ||
| 99 | + <td>是否使用均匀分布噪声(当前实现仅支持false)。默认值:true。</td> | ||
| 100 | + <td>Bool</td> | ||
| 101 | + <td>-</td> | ||
| 102 | + </tr> | ||
| 103 | + <tr> | ||
| 104 | + <td>noise</td> | ||
| 105 | + <td>属性</td> | ||
| 106 | + <td>噪声类型(当前实现仅支持"zero")。默认值:"uniform"。</td> | ||
| 107 | + <td>String</td> | ||
| 108 | + <td>-</td> | ||
| 109 | + </tr> | ||
| 110 | + </tbody></table> | ||
| 111 | + | ||
| 112 | +## 约束说明 | ||
| 113 | + | ||
| 114 | +- input 必须为 4D tensor。 | ||
| 115 | +- offsets 必须为 2D tensor,且第二维必须为 2。 | ||
| 116 | +- size 必须为 1D const tensor,且长度为 2。 | ||
| 117 | +- input 和 offsets 的 batch 维度必须一致。 | ||
| 118 | +- 当前仅支持 float32 数据类型。 | ||
| 119 | +- 当前仅支持 noise="zero" 的越界填充模式。 | ||
| 120 | + | ||
| 121 | +## 调用说明 | ||
| 122 | + | ||
| 123 | +| 调用方式 | 样例代码 | 说明 | | ||
| 124 | +| ---------------- | --------------------------- | --------------------------------------------------- | | ||
| 125 | +| 图模式 | - | 通过[算子IR](op_graph/extract_glimpse_v2_proto.h)构图方式调用ExtractGlimpseV2算子。 | | ||
| @@ -0,0 +1,40 @@ | |||
| 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 extract_glimpse_v2_graph_infer.cpp | ||
| 15 | + * \brief extract_glimpse_v2 operator graph infer resource | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +namespace ops { | ||
| 22 | +using namespace ge; | ||
| 23 | + | ||
| 24 | +static constexpr int64_t IDX_0 = 0; | ||
| 25 | + | ||
| 26 | +static ge::graphStatus InferDataTypeExtractGlimpseV2(gert::InferDataTypeContext* context) | ||
| 27 | +{ | ||
| 28 | + OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeExtractGlimpseV2"); | ||
| 29 | + | ||
| 30 | + // Output dtype is same as input dtype | ||
| 31 | + ge::DataType inputDtype = context->GetInputDataType(IDX_0); | ||
| 32 | + context->SetOutputDataType(IDX_0, inputDtype); | ||
| 33 | + | ||
| 34 | + OP_LOGD(context->GetNodeName(), "End to do InferDataTypeExtractGlimpseV2"); | ||
| 35 | + return GRAPH_SUCCESS; | ||
| 36 | +} | ||
| 37 | + | ||
| 38 | +IMPL_OP(ExtractGlimpseV2).InferDataType(InferDataTypeExtractGlimpseV2); | ||
| 39 | + | ||
| 40 | +}; // namespace ops | ||
| @@ -0,0 +1,56 @@ | |||
| 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 extract_glimpse_v2_proto.h | ||
| 15 | + * \brief Operator proto declaration for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +namespace ge { | ||
| 24 | + | ||
| 25 | +/** | ||
| 26 | + *@brief Extracts a glimpse of specified size from input images at given offsets. | ||
| 27 | + *@par Inputs: | ||
| 28 | + *Three inputs, including: | ||
| 29 | + * @li input: A 4D NHWC Tensor. Must be one of the following types: float32. | ||
| 30 | + * @li size: A 1D Tensor of shape [2]. Must be one of the following types: int32. | ||
| 31 | + * @li offsets: A 2D Tensor of shape [batch, 2]. Must be one of the following types: float32. \n | ||
| 32 | + | ||
| 33 | + *@par Outputs: | ||
| 34 | + *glimpse: A 4D NHWC Tensor of shape [batch, size_h, size_w, channels]. Must be one of the following types: float32. | ||
| 35 | + *@par Attributes: | ||
| 36 | + * @li centered: Bool. Whether offsets are centered. Default: true. | ||
| 37 | + * @li normalized: Bool. Whether offsets are normalized to [-1, 1]. Default: true. | ||
| 38 | + * @li uniform_noise: Bool. Whether to use uniform noise. Default: true. | ||
| 39 | + * @li noise: String. Noise type. Default: "uniform". | ||
| 40 | + *@par Third-party framework compatibility | ||
| 41 | + *Compatible with the TensorFlow operator ExtractGlimpseV2. | ||
| 42 | + */ | ||
| 43 | +REG_OP(ExtractGlimpseV2) | ||
| 44 | + .INPUT(input, TensorType({DT_FLOAT})) | ||
| 45 | + .INPUT(size, TensorType({DT_INT32})) | ||
| 46 | + .INPUT(offsets, TensorType({DT_FLOAT})) | ||
| 47 | + .OUTPUT(glimpse, TensorType({DT_FLOAT})) | ||
| 48 | + .ATTR(centered, Bool, true) | ||
| 49 | + .ATTR(normalized, Bool, true) | ||
| 50 | + .ATTR(uniform_noise, Bool, true) | ||
| 51 | + .ATTR(noise, String, "uniform") | ||
| 52 | + .OP_END_FACTORY_REG(ExtractGlimpseV2) | ||
| 53 | + | ||
| 54 | +} // namespace ge | ||
| 55 | + | ||
| 56 | + | ||
| @@ -0,0 +1,210 @@ | |||
| 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 extract_glimpse_v2_tiling.cpp | ||
| 15 | + * \brief Tiling implementation for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +namespace optiling { | ||
| 28 | + | ||
| 29 | +constexpr int64_t PER_CORE_MIN = 1024; | ||
| 30 | +constexpr uint32_t DCACHE_SIZE = 32 * 1024; | ||
| 31 | +constexpr uint32_t STATIC_UB_ESTIMATE = 0; | ||
| 32 | +constexpr uint32_t ATTR_IDX_CENTERED = 0; | ||
| 33 | +constexpr uint32_t ATTR_IDX_NORMALIZED = 1; | ||
| 34 | +constexpr uint32_t ATTR_IDX_UNIFORM_NOISE = 2; | ||
| 35 | +constexpr uint32_t ATTR_IDX_NOISE = 3; | ||
| 36 | +constexpr uint32_t SCENE_CENTERED_BIT = 2; | ||
| 37 | +constexpr uint32_t SCENE_NORMALIZED_BIT = 1; | ||
| 38 | +constexpr uint32_t WORKSPACE_COUNT = 1; | ||
| 39 | + | ||
| 40 | +struct ExtractGlimpseV2CompileInfo {}; | ||
| 41 | + | ||
| 42 | +struct ShapeInfo { | ||
| 43 | + int32_t batchSize; | ||
| 44 | + int32_t inputH; | ||
| 45 | + int32_t inputW; | ||
| 46 | + int32_t channels; | ||
| 47 | + int32_t glimpseH; | ||
| 48 | + int32_t glimpseW; | ||
| 49 | + int64_t maxCoreNum; | ||
| 50 | + uint64_t ubSize; | ||
| 51 | +}; | ||
| 52 | + | ||
| 53 | +static ge::graphStatus GetPlatformAndShapeInfo(gert::TilingContext* context, ShapeInfo& info) | ||
| 54 | +{ | ||
| 55 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 56 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 57 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 58 | + info.maxCoreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 59 | + OP_CHECK_IF(info.maxCoreNum == 0, OP_LOGE(context, "maxCoreNum is 0"), return ge::GRAPH_FAILED); | ||
| 60 | + info.ubSize = 0; | ||
| 61 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, info.ubSize); | ||
| 62 | + OP_CHECK_IF(info.ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 63 | + | ||
| 64 | + auto inputShape = context->GetInputShape(0); | ||
| 65 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputShape); | ||
| 66 | + int64_t batchSize = inputShape->GetShape().GetDim(0); | ||
| 67 | + int64_t inputH = inputShape->GetShape().GetDim(1); | ||
| 68 | + int64_t inputW = inputShape->GetShape().GetDim(2); | ||
| 69 | + int64_t channels = inputShape->GetShape().GetDim(3); | ||
| 70 | + | ||
| 71 | + auto outputShape = context->GetOutputShape(0); | ||
| 72 | + OP_CHECK_NULL_WITH_CONTEXT(context, outputShape); | ||
| 73 | + int64_t glimpseH = outputShape->GetShape().GetDim(1); | ||
| 74 | + int64_t glimpseW = outputShape->GetShape().GetDim(2); | ||
| 75 | + | ||
🟡 Medium |
只校验 问题链路:
对比
upsample_bicubic2d 有显式校验,extract_glimpse_v2 缺。 修复建议:
![]() ![]() | |||
| 76 | + if (glimpseH <= 0 || glimpseW <= 0) { | ||
| 77 | + std::string valMsg = "[" + std::to_string(glimpseH) + ", " + std::to_string(glimpseW) + "]"; | ||
| 78 | + OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", valMsg.c_str(), | ||
| 79 | + "output_size elements must be greater than 0"); | ||
| 80 | + return ge::GRAPH_FAILED; | ||
| 81 | + } | ||
| 82 | + | ||
| 83 | + OP_CHECK_IF(batchSize > INT32_MAX || inputH > INT32_MAX || inputW > INT32_MAX || channels > INT32_MAX || | ||
| 84 | + glimpseH > INT32_MAX || glimpseW > INT32_MAX, | ||
| 85 | + OP_LOGE(context, "dimension exceeds int32_t range"), return ge::GRAPH_FAILED); | ||
| 86 | + | ||
| 87 | + info.batchSize = static_cast<int32_t>(batchSize); | ||
| 88 | + info.inputH = static_cast<int32_t>(inputH); | ||
| 89 | + info.inputW = static_cast<int32_t>(inputW); | ||
| 90 | + info.channels = static_cast<int32_t>(channels); | ||
| 91 | + info.glimpseH = static_cast<int32_t>(glimpseH); | ||
| 92 | + info.glimpseW = static_cast<int32_t>(glimpseW); | ||
| 93 | + return ge::GRAPH_SUCCESS; | ||
| 94 | +} | ||
| 95 | + | ||
| 96 | +static ge::graphStatus ValidateAttributes(gert::TilingContext* context, bool& centered, bool& normalized) | ||
| 97 | +{ | ||
| 98 | + const gert::RuntimeAttrs* attrs = context->GetAttrs(); | ||
| 99 | + OP_CHECK_NULL_WITH_CONTEXT(context, attrs); | ||
| 100 | + const bool* centeredPtr = attrs->GetBool(ATTR_IDX_CENTERED); | ||
| 101 | + const bool* normalizedPtr = attrs->GetBool(ATTR_IDX_NORMALIZED); | ||
| 102 | + const bool* uniformNoisePtr = attrs->GetBool(ATTR_IDX_UNIFORM_NOISE); | ||
| 103 | + const char* noiseStr = attrs->GetStr(ATTR_IDX_NOISE); | ||
| 104 | + OP_CHECK_IF(centeredPtr == nullptr || normalizedPtr == nullptr || uniformNoisePtr == nullptr || noiseStr == nullptr, | ||
| 105 | + OP_LOGE(context, "attr pointer is null"), return ge::GRAPH_FAILED); | ||
| 106 | + if (*uniformNoisePtr) { | ||
| 107 | + OP_LOGE_FOR_INVALID_VALUE(context->GetNodeName(), "uniform_noise", "true", "false"); | ||
| 108 | + return ge::GRAPH_FAILED; | ||
| 109 | + } | ||
| 110 | + if (strcmp(noiseStr, "zero") != 0) { | ||
| 111 | + OP_LOGE_FOR_INVALID_VALUE(context->GetNodeName(), "noise", noiseStr, "zero"); | ||
| 112 | + return ge::GRAPH_FAILED; | ||
| 113 | + } | ||
| 114 | + centered = *centeredPtr; | ||
| 115 | + normalized = *normalizedPtr; | ||
| 116 | + return ge::GRAPH_SUCCESS; | ||
| 117 | +} | ||
| 118 | + | ||
| 119 | +static ge::graphStatus ComputeAndSetTiling(gert::TilingContext* context, const ShapeInfo& info, bool centered, | ||
| 120 | + bool normalized) | ||
| 121 | +{ | ||
| 122 | + int64_t outputTotalElements = static_cast<int64_t>(info.batchSize) * info.glimpseH * info.glimpseW * info.channels; | ||
| 123 | + | ||
| 124 | + int64_t perCoreElements = (outputTotalElements + info.maxCoreNum - 1) / info.maxCoreNum; | ||
| 125 | + if (perCoreElements < PER_CORE_MIN) { | ||
| 126 | + perCoreElements = PER_CORE_MIN; | ||
| 127 | + } | ||
| 128 | + int32_t needCoreNum = static_cast<int32_t>((outputTotalElements + perCoreElements - 1) / perCoreElements); | ||
| 129 | + if (needCoreNum > static_cast<int32_t>(info.maxCoreNum)) { | ||
| 130 | + needCoreNum = static_cast<int32_t>(info.maxCoreNum); | ||
| 131 | + } | ||
| 132 | + if (needCoreNum < 1) { | ||
| 133 | + needCoreNum = 1; | ||
| 134 | + } | ||
| 135 | + | ||
| 136 | + auto* tiling = context->GetTilingData<ExtractGlimpseV2TilingData>(); | ||
| 137 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 138 | + OP_CHECK_IF(memset_s(tiling, sizeof(ExtractGlimpseV2TilingData), 0, sizeof(ExtractGlimpseV2TilingData)) != EOK, | ||
| 139 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 140 | + tiling->needCoreNum = needCoreNum; | ||
| 141 | + tiling->batchSize = info.batchSize; | ||
| 142 | + tiling->inputH = info.inputH; | ||
| 143 | + tiling->inputW = info.inputW; | ||
| 144 | + tiling->channels = info.channels; | ||
| 145 | + tiling->glimpseH = info.glimpseH; | ||
| 146 | + tiling->glimpseW = info.glimpseW; | ||
| 147 | + tiling->centered = centered ? 1 : 0; | ||
| 148 | + tiling->normalized = normalized ? 1 : 0; | ||
| 149 | + tiling->outputTotalElements = outputTotalElements; | ||
| 150 | + tiling->perCoreElements = perCoreElements; | ||
| 151 | + | ||
| 152 | + context->SetBlockDim(needCoreNum); | ||
| 153 | + uint32_t sceneMode = (centered ? SCENE_CENTERED_BIT : 0) + (normalized ? SCENE_NORMALIZED_BIT : 0); | ||
| 154 | + context->SetTilingKey(GET_TPL_TILING_KEY(sceneMode)); | ||
| 155 | + return ge::GRAPH_SUCCESS; | ||
| 156 | +} | ||
| 157 | + | ||
| 158 | +static ge::graphStatus SetupMemoryAndWorkspace(gert::TilingContext* context, uint64_t ubSize) | ||
| 159 | +{ | ||
| 160 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 161 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 162 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 163 | + | ||
| 164 | + OP_CHECK_IF((ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE), | ||
| 165 | + OP_LOGE(context, "ubSize %lu <= DCACHE_SIZE + STATIC_UB_ESTIMATE", ubSize), return ge::GRAPH_FAILED); | ||
| 166 | + auto res = context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE)); | ||
| 167 | + OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context, "SetLocalMemorySize failed"), return ge::GRAPH_FAILED); | ||
| 168 | + | ||
| 169 | + uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize(); | ||
| 170 | + size_t* ws = context->GetWorkspaceSizes(WORKSPACE_COUNT); | ||
| 171 | + OP_CHECK_NULL_WITH_CONTEXT(context, ws); | ||
| 172 | + ws[0] = static_cast<size_t>(sysWorkspaceSize); | ||
| 173 | + return ge::GRAPH_SUCCESS; | ||
| 174 | +} | ||
| 175 | + | ||
| 176 | +static ge::graphStatus ExtractGlimpseV2TilingFunc(gert::TilingContext* context) | ||
| 177 | +{ | ||
| 178 | + ShapeInfo shapeInfo{}; | ||
| 179 | + if (GetPlatformAndShapeInfo(context, shapeInfo) != ge::GRAPH_SUCCESS) { | ||
| 180 | + return ge::GRAPH_FAILED; | ||
| 181 | + } | ||
| 182 | + | ||
| 183 | + bool centered = false; | ||
| 184 | + bool normalized = false; | ||
| 185 | + if (ValidateAttributes(context, centered, normalized) != ge::GRAPH_SUCCESS) { | ||
| 186 | + return ge::GRAPH_FAILED; | ||
| 187 | + } | ||
| 188 | + | ||
| 189 | + if (ComputeAndSetTiling(context, shapeInfo, centered, normalized) != ge::GRAPH_SUCCESS) { | ||
| 190 | + return ge::GRAPH_FAILED; | ||
| 191 | + } | ||
| 192 | + | ||
| 193 | + if (SetupMemoryAndWorkspace(context, shapeInfo.ubSize) != ge::GRAPH_SUCCESS) { | ||
| 194 | + return ge::GRAPH_FAILED; | ||
| 195 | + } | ||
| 196 | + | ||
| 197 | + return ge::GRAPH_SUCCESS; | ||
| 198 | +} | ||
| 199 | + | ||
| 200 | +static ge::graphStatus TilingParseForExtractGlimpseV2([[maybe_unused]] gert::TilingParseContext* context) | ||
| 201 | +{ | ||
| 202 | + return ge::GRAPH_SUCCESS; | ||
| 203 | +} | ||
| 204 | + | ||
| 205 | +IMPL_OP_OPTILING(ExtractGlimpseV2) | ||
| 206 | + .InputsDataDependency({1}) | ||
| 207 | + .Tiling(ExtractGlimpseV2TilingFunc) | ||
| 208 | + .TilingParse<ExtractGlimpseV2CompileInfo>(TilingParseForExtractGlimpseV2); | ||
| 209 | + | ||
| 210 | +} // namespace optiling | ||
| @@ -0,0 +1,13 @@ | |||
| 1 | +; 该文件主要影响 opc 工具 编译二进制kernel时, --simplified_key_mode 选项中填写的值,格式如下所示: | ||
| 2 | +; [某算子] | ||
| 3 | +; default=xx | ||
| 4 | +; ascendxx=xx | ||
| 5 | +; 其中,default为默认mode,ascendxx为可选mode,如果不同芯片有差异化要求时,需要配置; | ||
| 6 | +; 1)如果没有配置:非ascendC算子继续按空处理,即opc编译命令中不添加 --simplified_key_mode 选项,AscendC算子按照 simplified_key_mode=0 处理 | ||
| 7 | +; 2)如果仅有default配置:各个版本按default配置 | ||
| 8 | +; 3)如果仅有某些平台的配置,没有default配置:对应平台的按照配置的值传递,非对应平台的:非AscendC算子继续按空处理,AscendC算子按照 simplified_key_mode=0 处理 | ||
| 9 | +; 4)如果default配置和平台配置都有:对应平台的使用平台的配置,非对应的平台的以default值配置。 | ||
| 10 | +; 5)对于自定义simplified key的情况,需要在binary_simplified_key_mode.ini 文件中显式配置为None,不传入 --simplified_key_mode 选项,由opc工具和FE框架自行判断使用何种模式 | ||
| 11 | +; 6)是否是AscendC算子,由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定 | ||
| 12 | +[ExtractGlimpseV2] | ||
| 13 | +default=0 | ||
| @@ -0,0 +1,70 @@ | |||
| 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 extract_glimpse_v2_def.cpp | ||
| 15 | + * \brief Operator definition for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | +namespace ops { | ||
| 21 | + | ||
| 22 | +class ExtractGlimpseV2 : public OpDef { | ||
| 23 | +public: | ||
| 24 | + explicit ExtractGlimpseV2(const char* name) : OpDef(name) | ||
| 25 | + { | ||
| 26 | + this->Input("input") | ||
| 27 | + .ParamType(REQUIRED) | ||
| 28 | + .DataType({ge::DT_FLOAT}) | ||
| 29 | + .Format({ge::FORMAT_ND}) | ||
| 30 | + .UnknownShapeFormat({ge::FORMAT_ND}) | ||
| 31 | + .AutoContiguous(); | ||
| 32 | + this->Input("size") | ||
| 33 | + .ParamType(REQUIRED) | ||
| 34 | + .ValueDepend(OPTIONAL) | ||
| 35 | + .DataType({ge::DT_INT32}) | ||
| 36 | + .Format({ge::FORMAT_ND}) | ||
| 37 | + .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 38 | + this->Input("offsets") | ||
| 39 | + .ParamType(REQUIRED) | ||
| 40 | + .DataType({ge::DT_FLOAT}) | ||
| 41 | + .Format({ge::FORMAT_ND}) | ||
| 42 | + .UnknownShapeFormat({ge::FORMAT_ND}) | ||
| 43 | + .AutoContiguous(); | ||
| 44 | + this->Output("glimpse") | ||
| 45 | + .ParamType(REQUIRED) | ||
| 46 | + .DataType({ge::DT_FLOAT}) | ||
| 47 | + .Format({ge::FORMAT_ND}) | ||
| 48 | + .UnknownShapeFormat({ge::FORMAT_ND}) | ||
| 49 | + .AutoContiguous(); | ||
| 50 | + | ||
| 51 | + this->Attr("centered").AttrType(OPTIONAL).Bool(true); | ||
| 52 | + this->Attr("normalized").AttrType(OPTIONAL).Bool(true); | ||
| 53 | + this->Attr("uniform_noise").AttrType(OPTIONAL).Bool(true); | ||
| 54 | + this->Attr("noise").AttrType(OPTIONAL).String("uniform"); | ||
| 55 | + | ||
| 56 | + OpAICoreConfig aicoreConfig; | ||
| 57 | + aicoreConfig.DynamicCompileStaticFlag(true) | ||
| 58 | + .DynamicFormatFlag(false) | ||
| 59 | + .DynamicRankSupportFlag(true) | ||
| 60 | + .DynamicShapeSupportFlag(true) | ||
| 61 | + .NeedCheckSupportFlag(false) | ||
| 62 | + .PrecisionReduceFlag(false) | ||
| 63 | + .ExtendCfgInfo("opFile.value", "extract_glimpse_v2"); | ||
| 64 | + this->AICore().AddConfig("ascend950", aicoreConfig); | ||
| 65 | + } | ||
| 66 | +}; | ||
| 67 | + | ||
| 68 | +OP_ADD(ExtractGlimpseV2); | ||
| 69 | + | ||
| 70 | +} // namespace ops | ||
| @@ -0,0 +1,122 @@ | |||
| 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 extract_glimpse_v2_infershape.cpp | ||
| 15 | + * \brief InferShape implementation for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +using namespace ge; | ||
| 22 | + | ||
| 23 | +namespace ops { | ||
| 24 | + | ||
| 25 | +static constexpr int64_t IDX_0 = 0; | ||
| 26 | +static constexpr int64_t IDX_1 = 1; | ||
| 27 | +static constexpr int64_t IDX_2 = 2; | ||
| 28 | +static constexpr int64_t IDX_3 = 3; | ||
| 29 | +static constexpr size_t INPUT_REQUIRED_RANK = 4; | ||
| 30 | +static constexpr size_t OFFSETS_REQUIRED_RANK = 2; | ||
| 31 | +static constexpr int64_t OFFSETS_COORD_DIM = 2; | ||
| 32 | +static constexpr size_t SIZE_REQUIRED_RANK = 1; | ||
| 33 | +static constexpr int64_t SIZE_ELEMENT_COUNT = 2; | ||
| 34 | +static constexpr size_t OUTPUT_RANK = 4; | ||
| 35 | + | ||
| 36 | +static ge::graphStatus ValidateInputShapes(gert::InferShapeContext* context, const gert::Shape* inputShape, | ||
| 37 | + const gert::Shape* sizeShape, const gert::Shape* offsetsShape) | ||
| 38 | +{ | ||
| 39 | + if (inputShape->GetDimNum() != INPUT_REQUIRED_RANK) { | ||
| 40 | + OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "input", | ||
| 41 | + (std::to_string(inputShape->GetDimNum()) + "D").c_str(), | ||
| 42 | + (std::to_string(INPUT_REQUIRED_RANK) + "D").c_str()); | ||
| 43 | + return ge::GRAPH_FAILED; | ||
| 44 | + } | ||
| 45 | + if (offsetsShape->GetDimNum() != OFFSETS_REQUIRED_RANK) { | ||
| 46 | + OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "offsets", | ||
| 47 | + (std::to_string(offsetsShape->GetDimNum()) + "D").c_str(), | ||
| 48 | + (std::to_string(OFFSETS_REQUIRED_RANK) + "D").c_str()); | ||
| 49 | + return ge::GRAPH_FAILED; | ||
| 50 | + } | ||
| 51 | + if (offsetsShape->GetDim(1) != OFFSETS_COORD_DIM) { | ||
| 52 | + OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "offsets", std::to_string(offsetsShape->GetDim(1)).c_str(), | ||
| 53 | + std::to_string(OFFSETS_COORD_DIM).c_str()); | ||
| 54 | + return ge::GRAPH_FAILED; | ||
| 55 | + } | ||
| 56 | + if (inputShape->GetDim(0) != offsetsShape->GetDim(0)) { | ||
| 57 | + std::string dimMsg = std::to_string(inputShape->GetDim(0)) + " and " + std::to_string(offsetsShape->GetDim(0)); | ||
| 58 | + OP_LOGE_FOR_INVALID_SHAPEDIMS_WITH_REASON(context->GetNodeName(), "input and offsets", dimMsg.c_str(), | ||
| 59 | + "batch size of input and offsets must be equal"); | ||
| 60 | + return ge::GRAPH_FAILED; | ||
| 61 | + } | ||
| 62 | + if (sizeShape->GetDimNum() != SIZE_REQUIRED_RANK || sizeShape->GetDim(0) != SIZE_ELEMENT_COUNT) { | ||
| 63 | + OP_LOGE_FOR_INVALID_SHAPESIZE( | ||
| 64 | + context->GetNodeName(), "size", (std::to_string(sizeShape->GetDimNum()) + "D").c_str(), | ||
| 65 | + (std::to_string(SIZE_REQUIRED_RANK) + "D with " + std::to_string(SIZE_ELEMENT_COUNT) + " elements") | ||
| 66 | + .c_str()); | ||
| 67 | + return ge::GRAPH_FAILED; | ||
| 68 | + } | ||
| 69 | + return ge::GRAPH_SUCCESS; | ||
| 70 | +} | ||
| 71 | + | ||
| 72 | +static ge::graphStatus ReadSizeData(gert::InferShapeContext* context, int64_t& glimpseH, int64_t& glimpseW) | ||
| 73 | +{ | ||
| 74 | + const gert::Tensor* sizeTensor = context->GetInputTensor(IDX_1); | ||
| 75 | + OP_CHECK_NULL_WITH_CONTEXT(context, sizeTensor); | ||
| 76 | + const int32_t* sizeData = static_cast<const int32_t*>(sizeTensor->GetAddr()); | ||
| 77 | + OP_CHECK_IF(sizeData == nullptr, OP_LOGE(context, "size data is null"), return ge::GRAPH_FAILED); | ||
| 78 | + glimpseH = static_cast<int64_t>(sizeData[0]); | ||
| 79 | + glimpseW = static_cast<int64_t>(sizeData[1]); | ||
| 80 | + return ge::GRAPH_SUCCESS; | ||
| 81 | +} | ||
| 82 | + | ||
| 83 | +static ge::graphStatus SetOutputShapeAndDtype(gert::InferShapeContext* context, const gert::Shape* inputShape, | ||
| 84 | + int64_t glimpseH, int64_t glimpseW) | ||
| 85 | +{ | ||
| 86 | + gert::Shape* outShape = context->GetOutputShape(IDX_0); | ||
| 87 | + OP_CHECK_NULL_WITH_CONTEXT(context, outShape); | ||
| 88 | + outShape->SetDimNum(OUTPUT_RANK); | ||
| 89 | + outShape->SetDim(0, inputShape->GetDim(0)); | ||
| 90 | + outShape->SetDim(IDX_1, glimpseH); | ||
| 91 | + outShape->SetDim(IDX_2, glimpseW); | ||
| 92 | + outShape->SetDim(IDX_3, inputShape->GetDim(IDX_3)); | ||
| 93 | + auto* mutableOutDesc = context->GetComputeNodeInfo()->MutableOutputTdInfo(0); | ||
| 94 | + OP_CHECK_NULL_WITH_CONTEXT(context, mutableOutDesc); | ||
| 95 | + mutableOutDesc->SetDataType(context->GetInputDesc(0)->GetDataType()); | ||
| 96 | + return ge::GRAPH_SUCCESS; | ||
| 97 | +} | ||
| 98 | + | ||
| 99 | +static ge::graphStatus InferShapeExtractGlimpseV2(gert::InferShapeContext* context) | ||
| 100 | +{ | ||
| 101 | + const gert::Shape* inputShape = context->GetInputShape(IDX_0); | ||
| 102 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputShape); | ||
| 103 | + const gert::Shape* sizeShape = context->GetInputShape(IDX_1); | ||
| 104 | + OP_CHECK_NULL_WITH_CONTEXT(context, sizeShape); | ||
| 105 | + const gert::Shape* offsetsShape = context->GetInputShape(IDX_2); | ||
| 106 | + OP_CHECK_NULL_WITH_CONTEXT(context, offsetsShape); | ||
| 107 | + if (ValidateInputShapes(context, inputShape, sizeShape, offsetsShape) != ge::GRAPH_SUCCESS) { | ||
| 108 | + return ge::GRAPH_FAILED; | ||
| 109 | + } | ||
| 110 | + int64_t glimpseH = 0; | ||
| 111 | + int64_t glimpseW = 0; | ||
| 112 | + if (ReadSizeData(context, glimpseH, glimpseW) != ge::GRAPH_SUCCESS) { | ||
| 113 | + return ge::GRAPH_FAILED; | ||
| 114 | + } | ||
| 115 | + return SetOutputShapeAndDtype(context, inputShape, glimpseH, glimpseW); | ||
| 116 | +} | ||
| 117 | + | ||
| 118 | +IMPL_OP_INFERSHAPE(ExtractGlimpseV2) | ||
| 119 | + .InferShape(InferShapeExtractGlimpseV2) | ||
| 120 | + .InputsDataDependency({1}); | ||
| 121 | + | ||
| 122 | +} // namespace ops | ||
| @@ -0,0 +1,170 @@ | |||
| 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 extract_glimpse_v2_simt.h | ||
| 15 | + * \brief SIMT kernel implementation for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +namespace NsExtractGlimpseV2 { | ||
| 29 | + | ||
| 30 | +using namespace AscendC; | ||
| 31 | + | ||
| 32 | +constexpr uint32_t THREADS_UINT32 = 2048; | ||
| 33 | +constexpr uint32_t THREADS_UINT64 = 1024; | ||
| 34 | +constexpr size_t UINT32_BYTE_WIDTH = 4; | ||
| 35 | +constexpr int32_t OFFSETS_STRIDE = 2; | ||
| 36 | +constexpr float CENTERED_SCALE = 0.5f; | ||
| 37 | + | ||
| 38 | +template <typename IDX_T> | ||
| 39 | +static constexpr uint32_t THREADS = (sizeof(IDX_T) == UINT32_BYTE_WIDTH) ? THREADS_UINT32 : THREADS_UINT64; | ||
| 40 | + | ||
| 41 | +template <typename T, typename IDX_T> | ||
| 42 | +__simt_callee__ inline void ProcessPixel(IDX_T idx, int32_t b, int32_t oh, int32_t ow, int32_t c, int32_t inputW, | ||
| 43 | + int32_t channels, int32_t startY, int32_t startX, int32_t baseY, int32_t baseX, | ||
| 44 | + int32_t copyH, int32_t copyW, int64_t inputSliceSize, __gm__ T* input, | ||
| 45 | + __gm__ T* output) | ||
| 46 | +{ | ||
| 47 | + int32_t localY = oh - baseY; | ||
| 48 | + int32_t localX = ow - baseX; | ||
🔵 Low Priority
建议将 建议:将 startY/startX/endY/endX/baseY/baseX/copyH/copyW 的计算从 ProcessPixel 提升到 OpExtractGlimpseV2SimtKernel 中(在调用 ProcessPixel 之前),作为额外参数传入,避免每个输出像素重复计算相同的边界值。 ![]() ![]() | |||
| 49 | + if (copyH > 0 && copyW > 0 && localY >= 0 && localY < copyH && localX >= 0 && localX < copyW) { | ||
| 50 | + int32_t srcY = startY + localY; | ||
| 51 | + int32_t srcX = startX + localX; | ||
| 52 | + int64_t inputIdx = static_cast<int64_t>(b) * inputSliceSize + static_cast<int64_t>(srcY) * inputW * channels + | ||
| 53 | + static_cast<int64_t>(srcX) * channels + c; | ||
| 54 | + output[idx] = input[inputIdx]; | ||
| 55 | + } else { | ||
| 56 | + output[idx] = static_cast<T>(0); | ||
| 57 | + } | ||
| 58 | +} | ||
| 59 | + | ||
| 60 | +template <typename T, bool Centered, bool Normalized, typename IDX_T> | ||
| 61 | +__simt_vf__ __aicore__ __launch_bounds__(THREADS<IDX_T>) inline void OpExtractGlimpseV2SimtKernel( | ||
| 62 | + IDX_T outputTotalElements, int32_t inputH, int32_t inputW, int32_t channels, int32_t glimpseH, int32_t glimpseW, | ||
| 63 | + IDX_T magicChannels, IDX_T shiftChannels, IDX_T magicGlimpseW, IDX_T shiftGlimpseW, IDX_T magicGlimpseH, | ||
| 64 | + IDX_T shiftGlimpseH, __gm__ T* input, __gm__ float* offsets, __gm__ T* output) | ||
| 65 | +{ | ||
| 66 | + const int64_t inputSliceSize = static_cast<int64_t>(inputH) * inputW * channels; | ||
| 67 | + | ||
| 68 | + for (IDX_T idx = static_cast<IDX_T>(blockIdx.x * blockDim.x + threadIdx.x); idx < outputTotalElements; | ||
| 69 | + idx += static_cast<IDX_T>(blockDim.x * gridDim.x)) { | ||
| 70 | + IDX_T remaining = idx; | ||
| 71 | + | ||
🟢 Low |
这是设计选择,因为算子签名与 TensorFlow 算子签名一致(input, size, offsets, output)。kernel 只用 修复建议:
![]() ![]() | |||
| 72 | + IDX_T q1 = Simt::UintDiv<IDX_T>(remaining, magicChannels, shiftChannels); | ||
| 73 | + int32_t c = static_cast<int32_t>(remaining - q1 * static_cast<IDX_T>(channels)); | ||
| 74 | + remaining = q1; | ||
| 75 | + | ||
| 76 | + IDX_T q2 = Simt::UintDiv<IDX_T>(remaining, magicGlimpseW, shiftGlimpseW); | ||
| 77 | + int32_t ow = static_cast<int32_t>(remaining - q2 * static_cast<IDX_T>(glimpseW)); | ||
| 78 | + remaining = q2; | ||
| 79 | + | ||
| 80 | + IDX_T q3 = Simt::UintDiv<IDX_T>(remaining, magicGlimpseH, shiftGlimpseH); | ||
| 81 | + int32_t oh = static_cast<int32_t>(remaining - q3 * static_cast<IDX_T>(glimpseH)); | ||
| 82 | + int32_t b = static_cast<int32_t>(q3); | ||
| 83 | + | ||
| 84 | + float offsetY = offsets[b * OFFSETS_STRIDE]; | ||
| 85 | + float offsetX = offsets[b * OFFSETS_STRIDE + 1]; | ||
| 86 | + | ||
| 87 | + if (offsetY != offsetY || offsetX != offsetX) { | ||
| 88 | + output[idx] = static_cast<T>(0); | ||
| 89 | + continue; | ||
| 90 | + } | ||
| 91 | + | ||
| 92 | + if constexpr (Normalized) { | ||
| 93 | + offsetY = offsetY * static_cast<float>(inputH); | ||
| 94 | + offsetX = offsetX * static_cast<float>(inputW); | ||
| 95 | + } | ||
| 96 | + if constexpr (Centered) { | ||
| 97 | + if constexpr (Normalized) { | ||
| 98 | + offsetY = offsetY * CENTERED_SCALE + static_cast<float>(inputH) * CENTERED_SCALE - | ||
| 99 | + static_cast<float>(glimpseH) * CENTERED_SCALE; | ||
| 100 | + offsetX = offsetX * CENTERED_SCALE + static_cast<float>(inputW) * CENTERED_SCALE - | ||
| 101 | + static_cast<float>(glimpseW) * CENTERED_SCALE; | ||
| 102 | + } else { | ||
| 103 | + offsetY = offsetY + static_cast<float>(inputH) * CENTERED_SCALE; | ||
| 104 | + offsetX = offsetX + static_cast<float>(inputW) * CENTERED_SCALE; | ||
| 105 | + } | ||
| 106 | + } | ||
| 107 | + | ||
| 108 | + int32_t startY = static_cast<int32_t>(floorf(fmaxf(0.0f, offsetY))); | ||
| 109 | + int32_t startX = static_cast<int32_t>(floorf(fmaxf(0.0f, offsetX))); | ||
| 110 | + int32_t endY = static_cast<int32_t>( | ||
| 111 | + ceilf(fminf(static_cast<float>(inputH), offsetY + static_cast<float>(glimpseH)))); | ||
| 112 | + int32_t endX = static_cast<int32_t>( | ||
| 113 | + ceilf(fminf(static_cast<float>(inputW), offsetX + static_cast<float>(glimpseW)))); | ||
| 114 | + int32_t baseY = static_cast<int32_t>(floorf(fmaxf(0.0f, -offsetY))); | ||
| 115 | + int32_t baseX = static_cast<int32_t>(floorf(fmaxf(0.0f, -offsetX))); | ||
| 116 | + int32_t copyH = endY - startY; | ||
| 117 | + int32_t copyW = endX - startX; | ||
| 118 | + | ||
| 119 | + ProcessPixel<T, IDX_T>(idx, b, oh, ow, c, inputW, channels, startY, startX, baseY, baseX, copyH, copyW, | ||
| 120 | + inputSliceSize, input, output); | ||
| 121 | + } | ||
| 122 | +} | ||
| 123 | + | ||
| 124 | +template <typename T, bool Centered, bool Normalized> | ||
| 125 | +__aicore__ inline void Process(GM_ADDR input, GM_ADDR size, GM_ADDR offsets, GM_ADDR output, GM_ADDR workspace, | ||
| 126 | + const ExtractGlimpseV2TilingData* tilingData) | ||
| 127 | +{ | ||
| 128 | + int64_t outputTotalElements = tilingData->outputTotalElements; | ||
| 129 | + if (outputTotalElements == 0) { | ||
| 130 | + return; | ||
| 131 | + } | ||
| 132 | + | ||
| 133 | + __gm__ T* inputGm = (__gm__ T*)input; | ||
| 134 | + __gm__ float* offsetsGm = (__gm__ float*)offsets; | ||
| 135 | + __gm__ T* outputGm = (__gm__ T*)output; | ||
| 136 | + | ||
| 137 | + int32_t inputH = tilingData->inputH; | ||
| 138 | + int32_t inputW = tilingData->inputW; | ||
| 139 | + int32_t channels = tilingData->channels; | ||
| 140 | + int32_t glimpseH = tilingData->glimpseH; | ||
| 141 | + int32_t glimpseW = tilingData->glimpseW; | ||
| 142 | + | ||
| 143 | + if (outputTotalElements <= static_cast<int64_t>(INT32_MAX)) { | ||
| 144 | + uint32_t mCh = 0, sCh = 0; | ||
| 145 | + uint32_t mGW = 0, sGW = 0; | ||
| 146 | + uint32_t mGH = 0, sGH = 0; | ||
| 147 | + GetUintDivMagicAndShift<uint32_t>(mCh, sCh, static_cast<uint32_t>(channels)); | ||
| 148 | + GetUintDivMagicAndShift<uint32_t>(mGW, sGW, static_cast<uint32_t>(glimpseW)); | ||
| 149 | + GetUintDivMagicAndShift<uint32_t>(mGH, sGH, static_cast<uint32_t>(glimpseH)); | ||
| 150 | + | ||
| 151 | + asc_vf_call<OpExtractGlimpseV2SimtKernel<T, Centered, Normalized, uint32_t>>( | ||
| 152 | + dim3(THREADS<uint32_t>), static_cast<uint32_t>(outputTotalElements), inputH, inputW, channels, glimpseH, | ||
| 153 | + glimpseW, mCh, sCh, mGW, sGW, mGH, sGH, inputGm, offsetsGm, outputGm); | ||
| 154 | + } else { | ||
| 155 | + uint64_t mCh = 0, sCh = 0; | ||
| 156 | + uint64_t mGW = 0, sGW = 0; | ||
| 157 | + uint64_t mGH = 0, sGH = 0; | ||
| 158 | + GetUintDivMagicAndShift<uint64_t>(mCh, sCh, static_cast<uint64_t>(channels)); | ||
| 159 | + GetUintDivMagicAndShift<uint64_t>(mGW, sGW, static_cast<uint64_t>(glimpseW)); | ||
| 160 | + GetUintDivMagicAndShift<uint64_t>(mGH, sGH, static_cast<uint64_t>(glimpseH)); | ||
| 161 | + | ||
| 162 | + asc_vf_call<OpExtractGlimpseV2SimtKernel<T, Centered, Normalized, uint64_t>>( | ||
| 163 | + dim3(THREADS<uint64_t>), static_cast<uint64_t>(outputTotalElements), inputH, inputW, channels, glimpseH, | ||
| 164 | + glimpseW, mCh, sCh, mGW, sGW, mGH, sGH, inputGm, offsetsGm, outputGm); | ||
| 165 | + } | ||
| 166 | +} | ||
| 167 | + | ||
| 168 | +} // namespace NsExtractGlimpseV2 | ||
| 169 | + | ||
| 170 | + | ||
| @@ -0,0 +1,35 @@ | |||
| 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 extract_glimpse_v2_tiling_data.h | ||
| 15 | + * \brief tiling data struct for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +struct ExtractGlimpseV2TilingData { | ||
| 22 | + int32_t needCoreNum; | ||
| 23 | + int32_t batchSize; | ||
| 24 | + int32_t inputH; | ||
| 25 | + int32_t inputW; | ||
| 26 | + int32_t channels; | ||
| 27 | + int32_t glimpseH; | ||
| 28 | + int32_t glimpseW; | ||
| 29 | + int32_t centered; | ||
| 30 | + int32_t normalized; | ||
| 31 | + int64_t outputTotalElements; | ||
| 32 | + int64_t perCoreElements; | ||
| 33 | +}; | ||
| 34 | + | ||
| 35 | + | ||
| @@ -0,0 +1,37 @@ | |||
| 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 extract_glimpse_v2_tiling_key.h | ||
| 15 | + * \brief tiling key declaration for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +ASCENDC_TPL_ARGS_DECL( | ||
| 29 | + ExtractGlimpseV2, | ||
| 30 | + ASCENDC_TPL_UINT_DECL(schMode, 2, ASCENDC_TPL_UI_LIST, | ||
| 31 | + SCENE_MODE_0, SCENE_MODE_1, SCENE_MODE_2, SCENE_MODE_3)); | ||
| 32 | + | ||
| 33 | +ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL( | ||
| 34 | + ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, | ||
| 35 | + SCENE_MODE_0, SCENE_MODE_1, SCENE_MODE_2, SCENE_MODE_3))); | ||
| 36 | + | ||
| 37 | + | ||
| @@ -0,0 +1,41 @@ | |||
| 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 extract_glimpse_v2.cpp | ||
| 15 | + * \brief kernel entry for extract_glimpse_v2 | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | +template <uint32_t schMode> | ||
| 21 | +__global__ __aicore__ void extract_glimpse_v2( | ||
| 22 | + GM_ADDR input, GM_ADDR size, GM_ADDR offsets, | ||
| 23 | + GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling) | ||
| 24 | +{ | ||
| 25 | + REGISTER_TILING_DEFAULT(ExtractGlimpseV2TilingData); | ||
| 26 | + GET_TILING_DATA_WITH_STRUCT(ExtractGlimpseV2TilingData, tilingData, tiling); | ||
| 27 | + | ||
| 28 | + if constexpr (schMode == SCENE_MODE_0) { | ||
| 29 | + NsExtractGlimpseV2::Process<DTYPE_INPUT, false, false>( | ||
| 30 | + input, size, offsets, output, workspace, &tilingData); | ||
| 31 | + } else if constexpr (schMode == SCENE_MODE_1) { | ||
| 32 | + NsExtractGlimpseV2::Process<DTYPE_INPUT, false, true>( | ||
| 33 | + input, size, offsets, output, workspace, &tilingData); | ||
| 34 | + } else if constexpr (schMode == SCENE_MODE_2) { | ||
| 35 | + NsExtractGlimpseV2::Process<DTYPE_INPUT, true, false>( | ||
| 36 | + input, size, offsets, output, workspace, &tilingData); | ||
| 37 | + } else { | ||
| 38 | + NsExtractGlimpseV2::Process<DTYPE_INPUT, true, true>( | ||
| 39 | + input, size, offsets, output, workspace, &tilingData); | ||
| 40 | + } | ||
| 41 | +} | ||
| @@ -0,0 +1,18 @@ | |||
| 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 | +file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 14 | +foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 15 | + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 16 | + add_subdirectory(${SUB_DIR}) | ||
| 17 | + endif() | ||
| 18 | +endforeach() | ||
| @@ -0,0 +1,18 @@ | |||
| 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 | +file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 14 | +foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 15 | + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 16 | + add_subdirectory(${SUB_DIR}) | ||
| 17 | + endif() | ||
| 18 | +endforeach() | ||
| @@ -0,0 +1,25 @@ | |||
| 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 | +file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 14 | +if(UT_TEST_ALL OR OP_HOST_UT) | ||
| 15 | + if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950") | ||
| 16 | + add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35) | ||
| 17 | + endif() | ||
| 18 | + add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | ||
| 19 | +endif() | ||
| 20 | + | ||
| 21 | +foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 22 | + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 23 | + add_subdirectory(${SUB_DIR}) | ||
| 24 | + endif() | ||
| 25 | +endforeach() | ||
| @@ -0,0 +1,181 @@ | |||
| 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 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +using namespace std; | ||
| 20 | +using namespace ge; | ||
| 21 | + | ||
| 22 | +class ExtractGlimpseV2Tiling : public testing::Test { | ||
| 23 | +protected: | ||
| 24 | + static void SetUpTestCase() | ||
| 25 | + { | ||
| 26 | + std::cout << "ExtractGlimpseV2Tiling SetUp" << std::endl; | ||
| 27 | + } | ||
| 28 | + | ||
| 29 | + static void TearDownTestCase() | ||
| 30 | + { | ||
| 31 | + std::cout << "ExtractGlimpseV2Tiling TearDown" << std::endl; | ||
| 32 | + } | ||
| 33 | +}; | ||
| 34 | + | ||
| 35 | +static std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend950"}}; | ||
| 36 | + | ||
| 37 | +TEST_F(ExtractGlimpseV2Tiling, extract_glimpse_v2_0) | ||
| 38 | +{ | ||
| 39 | + struct ExtractGlimpseV2CompileInfo { | ||
| 40 | + } compileInfo; | ||
| 41 | + gert::TilingContextPara tilingContextPara( | ||
| 42 | + "ExtractGlimpseV2", | ||
| 43 | + { | ||
| 44 | + {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 45 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, | ||
| 46 | + {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 47 | + }, | ||
| 48 | + { | ||
| 49 | + {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 50 | + }, | ||
| 51 | + { | ||
| 52 | + gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 53 | + gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 54 | + gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 55 | + gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")), | ||
| 56 | + }, | ||
| 57 | + &compileInfo, | ||
| 58 | + "Ascend950", | ||
| 59 | + 64, | ||
| 60 | + 262144, | ||
| 61 | + 4096); | ||
| 62 | + uint64_t expectTilingKey = 3; | ||
| 63 | + string expectTilingData = "4294967297 12884901891 8589934595 4294967298 1 12 1024 "; | ||
| 64 | + std::vector<size_t> expectWorkspaces = {16777216}; | ||
| 65 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | ||
| 66 | +} | ||
| 67 | + | ||
| 68 | +TEST_F(ExtractGlimpseV2Tiling, extract_glimpse_v2_1) | ||
| 69 | +{ | ||
| 70 | + struct ExtractGlimpseV2CompileInfo { | ||
| 71 | + } compileInfo; | ||
| 72 | + gert::TilingContextPara tilingContextPara( | ||
| 73 | + "ExtractGlimpseV2", | ||
| 74 | + { | ||
| 75 | + {{{15, 13, 16, 16}, {15, 13, 16, 16}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 76 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, | ||
| 77 | + {{{15, 2}, {15, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 78 | + }, | ||
| 79 | + { | ||
| 80 | + {{{15, 2, 9, 16}, {15, 2, 9, 16}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 81 | + }, | ||
| 82 | + { | ||
| 83 | + gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 84 | + gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 85 | + gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 86 | + gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")), | ||
| 87 | + }, | ||
| 88 | + &compileInfo, | ||
| 89 | + "Ascend950", | ||
| 90 | + 64, | ||
| 91 | + 262144, | ||
| 92 | + 4096); | ||
| 93 | + uint64_t expectTilingKey = 0; | ||
| 94 | + string expectTilingData = "64424509445 68719476749 8589934608 9 0 4320 1024 "; | ||
| 95 | + std::vector<size_t> expectWorkspaces = {16777216}; | ||
| 96 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | ||
| 97 | +} | ||
| 98 | + | ||
| 99 | +TEST_F(ExtractGlimpseV2Tiling, fail_uniform_noise_true) | ||
| 100 | +{ | ||
| 101 | + struct ExtractGlimpseV2CompileInfo { | ||
| 102 | + } compileInfo; | ||
| 103 | + gert::TilingContextPara tilingContextPara( | ||
| 104 | + "ExtractGlimpseV2", | ||
| 105 | + { | ||
| 106 | + {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 107 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, | ||
| 108 | + {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 109 | + }, | ||
| 110 | + { | ||
| 111 | + {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 112 | + }, | ||
| 113 | + { | ||
| 114 | + gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 115 | + gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 116 | + gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 117 | + gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")), | ||
| 118 | + }, | ||
| 119 | + &compileInfo, | ||
| 120 | + "Ascend950", | ||
| 121 | + 64, | ||
| 122 | + 262144, | ||
| 123 | + 4096); | ||
| 124 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {}); | ||
| 125 | +} | ||
| 126 | + | ||
| 127 | +TEST_F(ExtractGlimpseV2Tiling, fail_noise_not_zero) | ||
| 128 | +{ | ||
| 129 | + struct ExtractGlimpseV2CompileInfo { | ||
| 130 | + } compileInfo; | ||
| 131 | + gert::TilingContextPara tilingContextPara( | ||
| 132 | + "ExtractGlimpseV2", | ||
| 133 | + { | ||
| 134 | + {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 135 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, | ||
| 136 | + {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 137 | + }, | ||
| 138 | + { | ||
| 139 | + {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 140 | + }, | ||
| 141 | + { | ||
| 142 | + gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 143 | + gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 144 | + gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 145 | + gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("uniform")), | ||
| 146 | + }, | ||
| 147 | + &compileInfo, | ||
| 148 | + "Ascend950", | ||
| 149 | + 64, | ||
| 150 | + 262144, | ||
| 151 | + 4096); | ||
| 152 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {}); | ||
| 153 | +} | ||
| 154 | + | ||
| 155 | +TEST_F(ExtractGlimpseV2Tiling, fail_noise_gaussian) | ||
| 156 | +{ | ||
| 157 | + struct ExtractGlimpseV2CompileInfo { | ||
| 158 | + } compileInfo; | ||
| 159 | + gert::TilingContextPara tilingContextPara( | ||
| 160 | + "ExtractGlimpseV2", | ||
| 161 | + { | ||
| 162 | + {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 163 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, | ||
| 164 | + {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 165 | + }, | ||
| 166 | + { | ||
| 167 | + {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 168 | + }, | ||
| 169 | + { | ||
| 170 | + gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 171 | + gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 172 | + gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 173 | + gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("gaussian")), | ||
| 174 | + }, | ||
| 175 | + &compileInfo, | ||
| 176 | + "Ascend950", | ||
| 177 | + 64, | ||
| 178 | + 262144, | ||
| 179 | + 4096); | ||
| 180 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {}); | ||
| 181 | +} | ||
| @@ -0,0 +1,68 @@ | |||
| 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 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +class ExtractGlimpseV2Infershape : public testing::Test | ||
| 19 | +{ | ||
| 20 | +protected: | ||
| 21 | + static void SetUpTestCase() | ||
| 22 | + { | ||
| 23 | + std::cout << "ExtractGlimpseV2Infershape SetUp" << std::endl; | ||
| 24 | + } | ||
| 25 | + | ||
| 26 | + static void TearDownTestCase() | ||
| 27 | + { | ||
| 28 | + std::cout << "ExtractGlimpseV2Infershape TearDown" << std::endl; | ||
| 29 | + } | ||
| 30 | +}; | ||
| 31 | + | ||
| 32 | +TEST_F(ExtractGlimpseV2Infershape, extract_glimpse_v2_infershape_test1) | ||
| 33 | +{ | ||
| 34 | + int32_t sizeData1[2] = {2, 2}; | ||
| 35 | + gert::InfershapeContextPara infershapeContextPara( | ||
| 36 | + "ExtractGlimpseV2", | ||
| 37 | + { | ||
| 38 | + {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 39 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, sizeData1}, | ||
| 40 | + {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 41 | + }, | ||
| 42 | + { | ||
| 43 | + {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 44 | + }); | ||
| 45 | + std::vector<std::vector<int64_t>> expectOutputShape = { | ||
| 46 | + {1, 2, 2, 3}, | ||
| 47 | + }; | ||
| 48 | + ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape); | ||
| 49 | +} | ||
| 50 | + | ||
| 51 | +TEST_F(ExtractGlimpseV2Infershape, extract_glimpse_v2_infershape_test2) | ||
| 52 | +{ | ||
| 53 | + int32_t sizeData2[2] = {2, 9}; | ||
| 54 | + gert::InfershapeContextPara infershapeContextPara( | ||
| 55 | + "ExtractGlimpseV2", | ||
| 56 | + { | ||
| 57 | + {{{15, 13, 16, 16}, {15, 13, 16, 16}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 58 | + {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, sizeData2}, | ||
| 59 | + {{{15, 2}, {15, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 60 | + }, | ||
| 61 | + { | ||
| 62 | + {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}, | ||
| 63 | + }); | ||
| 64 | + std::vector<std::vector<int64_t>> expectOutputShape = { | ||
| 65 | + {15, 2, 9, 16}, | ||
| 66 | + }; | ||
| 67 | + ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape); | ||
| 68 | +} | ||
| @@ -4,7 +4,7 @@ | |||
| 4 | 4 | ||
| 5 | |产品 | 是否支持 | | 5 | |产品 | 是否支持 | |
| 6 | |:-------------------------|:----------:| | 6 | |:-------------------------|:----------:| |
| 7 | -| <term>Ascend 950PR/Ascend 950DT</term> | × | | 7 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | |
| 8 | | <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | | 8 | | <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | |
| 9 | | <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | | 9 | | <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | |
| 10 | | <term>Atlas 200I/500 A2 推理产品</term> | √ | | 10 | | <term>Atlas 200I/500 A2 推理产品</term> | √ | |
| @@ -8,6 +8,8 @@ | |||
| 8 | * 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. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 11 | /*! | 13 | /*! |
| 12 | * \file upsample_bicubic2d_def.cpp | 14 | * \file upsample_bicubic2d_def.cpp |
| 13 | * \brief | 15 | * \brief |
| @@ -38,6 +40,15 @@ public: | |||
| 38 | this->AICore().AddConfig("ascend910b"); | 40 | this->AICore().AddConfig("ascend910b"); |
| 39 | this->AICore().AddConfig("ascend910_93"); | 41 | this->AICore().AddConfig("ascend910_93"); |
| 40 | 42 | ||
| 43 | + OpAICoreConfig aicoreConfig; | ||
| 44 | + aicoreConfig.DynamicCompileStaticFlag(true) | ||
| 45 | + .DynamicRankSupportFlag(true) | ||
| 46 | + .DynamicShapeSupportFlag(true) | ||
| 47 | + .DynamicFormatFlag(true) | ||
| 48 | + .NeedCheckSupportFlag(false) | ||
| 49 | + .ExtendCfgInfo("opFile.value", "upsample_bicubic2d_apt"); | ||
| 50 | + this->AICore().AddConfig("ascend950", aicoreConfig); | ||
| 51 | + | ||
| 41 | OpAICoreConfig config310p; | 52 | OpAICoreConfig config310p; |
| 42 | config310p.DynamicCompileStaticFlag(true) | 53 | config310p.DynamicCompileStaticFlag(true) |
| 43 | .DynamicFormatFlag(true) | 54 | .DynamicFormatFlag(true) |
| @@ -8,6 +8,8 @@ | |||
| 8 | * 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. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| 11 | +/* Generated By CANNBot */ | ||
| 12 | + | ||
| 11 | /*! | 13 | /*! |
| 12 | * \file upsample_bicubic2d_tiling.cpp | 14 | * \file upsample_bicubic2d_tiling.cpp |
| 13 | * \brief | 15 | * \brief |
| @@ -18,7 +20,11 @@ | |||
| 18 | 20 | ||
| 19 | 21 | ||
| 20 | 22 | ||
| 21 | -#include "op_host/tiling_base.h" | 23 | +#include "op_host/tiling_util.h" |
| 24 | + | ||
| 25 | +namespace optiling { | ||
| 26 | +ge::graphStatus UpsampleBicubic2dTilingFunc(gert::TilingContext* context); | ||
| 27 | +} | ||
| 22 | 28 | ||
| 23 | namespace optiling { | 29 | namespace optiling { |
| 24 | constexpr uint32_t BEST_PERFORMANCE_SIZE_1 = 16; | 30 | constexpr uint32_t BEST_PERFORMANCE_SIZE_1 = 16; |
| @@ -583,6 +589,10 @@ void UpsampleBicubic2dTiling::FillTilingData() | |||
| 583 | 589 | ||
| 584 | static ge::graphStatus tiling4UpsampleBicubic2dTiling(gert::TilingContext* context) | 590 | static ge::graphStatus tiling4UpsampleBicubic2dTiling(gert::TilingContext* context) |
| 585 | { | 591 | { |
| 592 | + bool regBase = Ops::Cv::OpTiling::IsRegbaseSocVersion(context); | ||
| 593 | + if (regBase) { | ||
| 594 | + return optiling::UpsampleBicubic2dTilingFunc(context); | ||
| 595 | + } | ||
| 586 | UpsampleBicubic2dTiling tilingObject(context); | 596 | UpsampleBicubic2dTiling tilingObject(context); |
| 587 | context->SetScheduleMode(SCHEDULE_MODE); | 597 | context->SetScheduleMode(SCHEDULE_MODE); |
| 588 | return tilingObject.RunBigKernelTiling(); | 598 | return tilingObject.RunBigKernelTiling(); |
| @@ -0,0 +1,228 @@ | |||
| 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 upsample_bicubic2d_tiling.cpp | ||
| 15 | + * \brief tiling implementation for upsample_bicubic2d | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +namespace optiling { | ||
| 28 | + | ||
| 29 | +using namespace Ops::Cv::OpTiling; | ||
| 30 | + | ||
| 31 | +constexpr uint32_t DCACHE_SIZE = 128 * 1024; | ||
| 32 | +constexpr uint32_t STATIC_UB_ESTIMATE = 0; | ||
| 33 | +constexpr int64_t PER_CORE_MIN = 1024; | ||
| 34 | +constexpr int32_t INPUT_DIM_NUM = 4; | ||
| 35 | +constexpr int64_t PER_CORE_ALIGN = 32; | ||
| 36 | +constexpr int64_t ALIGN_MASK = PER_CORE_ALIGN - 1; | ||
| 37 | +constexpr int32_t OUTPUT_SIZE_ATTR_COUNT = 2; | ||
| 38 | +constexpr uint32_t ATTR_IDX_OUTPUT_SIZE = 0; | ||
| 39 | +constexpr uint32_t ATTR_IDX_ALIGN_CORNERS = 1; | ||
| 40 | +constexpr uint32_t ATTR_IDX_SCALES_H = 2; | ||
| 41 | +constexpr uint32_t ATTR_IDX_SCALES_W = 3; | ||
| 42 | + | ||
| 43 | +struct UpsampleBicubic2dCompileInfo {}; | ||
| 44 | + | ||
| 45 | +// 获取平台信息 | ||
| 46 | +static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum) | ||
| 47 | +{ | ||
| 48 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 49 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 50 | + | ||
| 51 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 52 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize); | ||
| 53 | + OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 54 | + | ||
| 55 | + coreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 56 | + OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED); | ||
| 57 | + | ||
| 58 | + return ge::GRAPH_SUCCESS; | ||
| 59 | +} | ||
| 60 | + | ||
| 61 | +// 获取shape和属性信息 | ||
| 62 | +static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int32_t& batchSize, int32_t& channels, | ||
| 63 | + int32_t& inputHeight, int32_t& inputWidth, int32_t& outputHeight, | ||
| 64 | + int32_t& outputWidth, bool& alignCorners, float& scalesH, float& scalesW) | ||
| 65 | +{ | ||
| 66 | + // 获取输入shape | ||
| 67 | + auto inputX = context->GetInputShape(0); | ||
| 68 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputX); | ||
| 69 | + auto inputShapeX = inputX->GetStorageShape(); | ||
| 70 | + | ||
| 71 | + if (inputShapeX.GetDimNum() != INPUT_DIM_NUM) { | ||
| 72 | + OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "x", | ||
| 73 | + (std::to_string(inputShapeX.GetDimNum()) + "D").c_str(), | ||
| 74 | + (std::to_string(INPUT_DIM_NUM) + "D").c_str()); | ||
| 75 | + return ge::GRAPH_FAILED; | ||
| 76 | + } | ||
| 77 | + | ||
| 78 | + batchSize = static_cast<int32_t>(inputShapeX.GetDim(0)); | ||
| 79 | + channels = static_cast<int32_t>(inputShapeX.GetDim(1)); | ||
| 80 | + inputHeight = static_cast<int32_t>(inputShapeX.GetDim(2)); | ||
| 81 | + inputWidth = static_cast<int32_t>(inputShapeX.GetDim(3)); | ||
| 82 | + | ||
| 83 | + // 获取属性 | ||
| 84 | + const gert::RuntimeAttrs* attrs = context->GetAttrs(); | ||
| 85 | + OP_CHECK_NULL_WITH_CONTEXT(context, attrs); | ||
| 86 | + | ||
| 87 | + // output_size (REQUIRED_ATTR, ListInt, index 0) | ||
| 88 | + const auto* outputSizeAttr = attrs->GetListInt(ATTR_IDX_OUTPUT_SIZE); | ||
| 89 | + OP_CHECK_NULL_WITH_CONTEXT(context, outputSizeAttr); | ||
| 90 | + if (outputSizeAttr->GetSize() != OUTPUT_SIZE_ATTR_COUNT) { | ||
| 91 | + OP_LOGE_FOR_INVALID_LISTSIZE(context->GetNodeName(), "output_size", | ||
| 92 | + std::to_string(outputSizeAttr->GetSize()).c_str(), | ||
| 93 | + std::to_string(OUTPUT_SIZE_ATTR_COUNT).c_str()); | ||
| 94 | + return ge::GRAPH_FAILED; | ||
| 95 | + } | ||
| 96 | + const int64_t* outputSizeData = outputSizeAttr->GetData(); | ||
| 97 | + outputHeight = static_cast<int32_t>(outputSizeData[0]); | ||
| 98 | + outputWidth = static_cast<int32_t>(outputSizeData[1]); | ||
| 99 | + if (outputHeight <= 0 || outputWidth <= 0) { | ||
| 100 | + std::string valMsg = "[" + std::to_string(outputHeight) + ", " + std::to_string(outputWidth) + "]"; | ||
| 101 | + OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", valMsg.c_str(), | ||
| 102 | + "output_size elements must be greater than 0"); | ||
| 103 | + return ge::GRAPH_FAILED; | ||
| 104 | + } | ||
| 105 | + | ||
| 106 | + // align_corners (ATTR, Bool, index 1) | ||
| 107 | + const bool* alignCornersPtr = attrs->GetBool(ATTR_IDX_ALIGN_CORNERS); | ||
| 108 | + OP_CHECK_NULL_WITH_CONTEXT(context, alignCornersPtr); | ||
| 109 | + alignCorners = *alignCornersPtr; | ||
| 110 | + | ||
| 111 | + // scales_h (ATTR, Float, index 2) | ||
| 112 | + const float* scalesHPtr = attrs->GetFloat(ATTR_IDX_SCALES_H); | ||
| 113 | + OP_CHECK_NULL_WITH_CONTEXT(context, scalesHPtr); | ||
| 114 | + scalesH = *scalesHPtr; | ||
| 115 | + | ||
| 116 | + // scales_w (ATTR, Float, index 3) | ||
| 117 | + const float* scalesWPtr = attrs->GetFloat(ATTR_IDX_SCALES_W); | ||
| 118 | + OP_CHECK_NULL_WITH_CONTEXT(context, scalesWPtr); | ||
| 119 | + scalesW = *scalesWPtr; | ||
| 120 | + | ||
| 121 | + return ge::GRAPH_SUCCESS; | ||
| 122 | +} | ||
| 123 | + | ||
| 124 | +static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context) | ||
| 125 | +{ | ||
| 126 | + int64_t userWorkspaceSize = 0; | ||
| 127 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo()); | ||
| 128 | + uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize(); | ||
| 129 | + size_t* currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 130 | + OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace); | ||
| 131 | + currentWorkspace[0] = static_cast<size_t>(userWorkspaceSize + static_cast<int64_t>(sysWorkspaceSize)); | ||
| 132 | + return ge::GRAPH_SUCCESS; | ||
| 133 | +} | ||
| 134 | + | ||
| 135 | +static void ComputeScaleFactors(int32_t inputHeight, int32_t inputWidth, int32_t outputHeight, int32_t outputWidth, | ||
| 136 | + bool alignCorners, float& actualScaleH, float& actualScaleW) | ||
| 137 | +{ | ||
| 138 | + if (alignCorners) { | ||
| 139 | + actualScaleH = (outputHeight > 1) ? static_cast<float>(inputHeight - 1) / static_cast<float>(outputHeight - 1) : | ||
| 140 | + 0.0f; | ||
| 141 | + actualScaleW = (outputWidth > 1) ? static_cast<float>(inputWidth - 1) / static_cast<float>(outputWidth - 1) : | ||
| 142 | + 0.0f; | ||
| 143 | + } else { | ||
| 144 | + actualScaleH = static_cast<float>(inputHeight) / static_cast<float>(outputHeight); | ||
| 145 | + actualScaleW = static_cast<float>(inputWidth) / static_cast<float>(outputWidth); | ||
| 146 | + } | ||
| 147 | +} | ||
| 148 | + | ||
| 149 | +static ge::graphStatus SetTilingAndCoreSplit(gert::TilingContext* context, int32_t batchSize, int32_t channels, | ||
| 150 | + int32_t inputHeight, int32_t inputWidth, int32_t outputHeight, | ||
| 151 | + int32_t outputWidth, float actualScaleH, float actualScaleW, | ||
| 152 | + int64_t coreNum) | ||
| 153 | +{ | ||
| 154 | + UpsampleBicubic2dRegBaseTilingData* tiling = context->GetTilingData<UpsampleBicubic2dRegBaseTilingData>(); | ||
| 155 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 156 | + OP_CHECK_IF(memset_s(tiling, sizeof(UpsampleBicubic2dRegBaseTilingData), 0, | ||
| 157 | + sizeof(UpsampleBicubic2dRegBaseTilingData)) != EOK, | ||
| 158 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 159 | + | ||
| 160 | + int64_t totalElements = static_cast<int64_t>(batchSize) * channels * outputHeight * outputWidth; | ||
| 161 | + tiling->totalElements = totalElements; | ||
| 162 | + | ||
| 163 | + int64_t perCoreElements = Ops::Base::CeilDiv(totalElements, coreNum); | ||
| 164 | + if (perCoreElements < PER_CORE_MIN) { | ||
| 165 | + perCoreElements = PER_CORE_MIN; | ||
| 166 | + } | ||
| 167 | + perCoreElements = ((perCoreElements + ALIGN_MASK) / PER_CORE_ALIGN) * PER_CORE_ALIGN; | ||
| 168 | + int64_t needCoreNum = Ops::Base::CeilDiv(totalElements, perCoreElements); | ||
| 169 | + if (needCoreNum < 1) { | ||
| 170 | + needCoreNum = 1; | ||
| 171 | + } | ||
| 172 | + tiling->needCoreNum = static_cast<int32_t>(needCoreNum); | ||
| 173 | + context->SetBlockDim(static_cast<uint32_t>(needCoreNum)); | ||
| 174 | + tiling->ncSize = batchSize * channels; | ||
| 175 | + tiling->inputHeight = inputHeight; | ||
| 176 | + tiling->inputWidth = inputWidth; | ||
| 177 | + tiling->outputHeight = outputHeight; | ||
| 178 | + tiling->outputWidth = outputWidth; | ||
| 179 | + tiling->scaleH = actualScaleH; | ||
| 180 | + tiling->scaleW = actualScaleW; | ||
| 181 | + return ge::GRAPH_SUCCESS; | ||
| 182 | +} | ||
| 183 | + | ||
| 184 | +static ge::graphStatus SetMemoryAndTilingKey(gert::TilingContext* context, uint64_t ubSize, bool alignCorners) | ||
| 185 | +{ | ||
| 186 | + OP_CHECK_IF((ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE), | ||
| 187 | + OP_LOGE(context, "ubSize %lu <= DCACHE_SIZE + STATIC_UB_ESTIMATE", ubSize), return ge::GRAPH_FAILED); | ||
| 188 | + auto res = context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE)); | ||
| 189 | + OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context, "SetLocalMemorySize failed"), return ge::GRAPH_FAILED); | ||
| 190 | + | ||
| 191 | + uint64_t tilingKey = GET_TPL_TILING_KEY(static_cast<uint64_t>( | ||
| 192 | + alignCorners ? UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE : UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE)); | ||
| 193 | + context->SetTilingKey(tilingKey); | ||
| 194 | + return ge::GRAPH_SUCCESS; | ||
🟢 Low |
实际影响:无功能影响,但代码可读性差。 修复建议:
![]() ![]() | |||
| 195 | +} | ||
| 196 | + | ||
| 197 | +ge::graphStatus UpsampleBicubic2dTilingFunc(gert::TilingContext* context) | ||
| 198 | +{ | ||
| 199 | + uint64_t ubSize; | ||
| 200 | + int64_t coreNum; | ||
| 201 | + OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, | ||
| 202 | + OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED); | ||
| 203 | + | ||
| 204 | + int32_t batchSize, channels, inputHeight, inputWidth; | ||
| 205 | + int32_t outputHeight, outputWidth; | ||
| 206 | + bool alignCorners; | ||
| 207 | + float scalesH, scalesW; | ||
| 208 | + OP_CHECK_IF(GetShapeAttrsInfo(context, batchSize, channels, inputHeight, inputWidth, outputHeight, outputWidth, | ||
| 209 | + alignCorners, scalesH, scalesW) != ge::GRAPH_SUCCESS, | ||
| 210 | + OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED); | ||
| 211 | + | ||
| 212 | + OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"), | ||
| 213 | + return ge::GRAPH_FAILED); | ||
| 214 | + | ||
| 215 | + float actualScaleH, actualScaleW; | ||
| 216 | + ComputeScaleFactors(inputHeight, inputWidth, outputHeight, outputWidth, alignCorners, actualScaleH, actualScaleW); | ||
| 217 | + | ||
| 218 | + OP_CHECK_IF(SetTilingAndCoreSplit(context, batchSize, channels, inputHeight, inputWidth, outputHeight, outputWidth, | ||
| 219 | + actualScaleH, actualScaleW, coreNum) != ge::GRAPH_SUCCESS, | ||
| 220 | + OP_LOGE(context, "SetTilingAndCoreSplit error"), return ge::GRAPH_FAILED); | ||
| 221 | + | ||
| 222 | + OP_CHECK_IF(SetMemoryAndTilingKey(context, ubSize, alignCorners) != ge::GRAPH_SUCCESS, | ||
| 223 | + OP_LOGE(context, "SetMemoryAndTilingKey error"), return ge::GRAPH_FAILED); | ||
| 224 | + | ||
| 225 | + return ge::GRAPH_SUCCESS; | ||
| 226 | +} | ||
| 227 | + | ||
| 228 | +} // namespace optiling | ||
| @@ -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 upsample_bicubic2d_simt.h | ||
| 15 | + * \brief SIMT kernel implementation for upsample_bicubic2d | ||
| 16 | + * | ||
| 17 | + * 性能优化(Stage 1): | ||
| 18 | + * R001/R007: AlignCorners 模板参数化,消除循环内运行期分支 | ||
| 19 | + * R002: 3D 坐标拆解使用 Simt::UintDiv 快除替代原生除法 | ||
| 20 | + * R003: 索引类型模板化(uint32/uint64),32 位路径吞吐翻倍 | ||
| 21 | + * R005: N*C 合轴降为 3D 拆解,减少 1 次 UintDiv | ||
| 22 | + */ | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + | ||
| 39 | +namespace NsUpsampleBicubic2d { | ||
| 40 | + | ||
| 41 | +using namespace AscendC; | ||
| 42 | + | ||
| 43 | +constexpr uint32_t BICUBIC_THREADS_32 = 512; | ||
| 44 | +constexpr uint32_t BICUBIC_THREADS_64 = 256; | ||
| 45 | +constexpr size_t UINT32_BYTE_SIZE = 4; | ||
| 46 | +constexpr int32_t BICUBIC_WINDOW_SIZE = 4; | ||
| 47 | +constexpr float CUBIC_A = -0.75f; | ||
| 48 | +constexpr float CUBIC_C2_COEFF_5A = 5.0f; | ||
| 49 | +constexpr float CUBIC_C2_COEFF_8A = 8.0f; | ||
| 50 | +constexpr float CUBIC_C2_COEFF_4A = 4.0f; | ||
| 51 | +constexpr float CUBIC_C1_COEFF_2 = 2.0f; | ||
| 52 | +constexpr float CUBIC_C1_COEFF_3 = 3.0f; | ||
| 53 | +constexpr float CUBIC_OFFSET_1 = 1.0f; | ||
| 54 | +constexpr float CUBIC_OFFSET_2 = 2.0f; | ||
| 55 | +constexpr float HALF_PIXEL_OFFSET = 0.5f; | ||
| 56 | + | ||
| 57 | +// ===== 线程数按索引位宽分级(R003/R006)===== | ||
| 58 | +// bicubic 插值寄存器压力中等偏高(16 像素读取 + 5 次三次插值) | ||
| 59 | +// uint32 路径: 512 线程;uint64 路径: 256 线程 | ||
| 60 | +template <typename UintT> | ||
| 61 | +static constexpr uint32_t BICUBIC_THREADS = (sizeof(UintT) == UINT32_BYTE_SIZE) ? BICUBIC_THREADS_32 : BICUBIC_THREADS_64; | ||
| 62 | + | ||
| 63 | +// ===== 类型转换辅助函数(__simt_callee__)===== | ||
| 64 | + | ||
| 65 | +template <typename T> | ||
| 66 | +__simt_callee__ __aicore__ inline float ToFloat(T v) | ||
| 67 | +{ | ||
| 68 | + if constexpr (std::is_same_v<T, float>) { | ||
| 69 | + return v; | ||
| 70 | + } else if constexpr (std::is_same_v<T, half>) { | ||
| 71 | + return __half2float(v); | ||
| 72 | + } else { | ||
| 73 | + return __bfloat162float(v); | ||
| 74 | + } | ||
| 75 | +} | ||
| 76 | + | ||
| 77 | +template <typename T> | ||
| 78 | +__simt_callee__ __aicore__ inline T FromFloat(float v) | ||
| 79 | +{ | ||
| 80 | + if constexpr (std::is_same_v<T, float>) { | ||
| 81 | + return v; | ||
| 82 | + } else if constexpr (std::is_same_v<T, half>) { | ||
| 83 | + return __float2half(v); | ||
| 84 | + } else { | ||
| 85 | + return __float2bfloat16(v); | ||
| 86 | + } | ||
| 87 | +} | ||
| 88 | + | ||
| 89 | +// ===== 三次插值核心函数(__simt_callee__)===== | ||
| 90 | + | ||
| 91 | +__simt_callee__ __aicore__ inline float CubicInterp1d( | ||
| 92 | + float x0, float x1, float x2, float x3, float t) | ||
| 93 | +{ | ||
| 94 | + const float A = CUBIC_A; | ||
| 95 | + // coeffs[0]: cubic_convolution2(t + 1, A) | ||
| 96 | + float x_0 = t + CUBIC_OFFSET_1; | ||
| 97 | + float c0 = ((A * x_0 - CUBIC_C2_COEFF_5A * A) * x_0 + CUBIC_C2_COEFF_8A * A) * x_0 - CUBIC_C2_COEFF_4A * A; | ||
| 98 | + // coeffs[1]: cubic_convolution1(t, A) | ||
| 99 | + float c1 = ((A + CUBIC_C1_COEFF_2) * t - (A + CUBIC_C1_COEFF_3)) * t * t + CUBIC_OFFSET_1; | ||
| 100 | + // coeffs[2]: cubic_convolution1(1 - t, A) | ||
| 101 | + float x2v = CUBIC_OFFSET_1 - t; | ||
| 102 | + float c2 = ((A + CUBIC_C1_COEFF_2) * x2v - (A + CUBIC_C1_COEFF_3)) * x2v * x2v + CUBIC_OFFSET_1; | ||
| 103 | + // coeffs[3]: cubic_convolution2(2 - t, A) | ||
| 104 | + float x_3 = CUBIC_OFFSET_2 - t; | ||
| 105 | + float c3 = ((A * x_3 - CUBIC_C2_COEFF_5A * A) * x_3 + CUBIC_C2_COEFF_8A * A) * x_3 - CUBIC_C2_COEFF_4A * A; | ||
| 106 | + return x0 * c0 + x1 * c1 + x2 * c2 + x3 * c3; | ||
| 107 | +} | ||
| 108 | + | ||
| 109 | +// ===== 边界安全像素读取(__simt_callee__)===== | ||
| 110 | + | ||
| 111 | +template <typename T> | ||
| 112 | +__simt_callee__ __aicore__ inline float GetPixelBounded( | ||
| 113 | + __gm__ T* data, int64_t ncBase, | ||
| 114 | + int32_t height, int32_t width, | ||
| 115 | + int32_t y, int32_t x) | ||
| 116 | +{ | ||
| 117 | + int32_t accessY = (y < 0) ? 0 : ((y >= height) ? (height - 1) : y); | ||
| 118 | + int32_t accessX = (x < 0) ? 0 : ((x >= width) ? (width - 1) : x); | ||
| 119 | + int64_t pixelIdx = ncBase + static_cast<int64_t>(accessY) * width + accessX; | ||
| 120 | + return ToFloat<T>(data[pixelIdx]); | ||
| 121 | +} | ||
| 122 | + | ||
| 123 | +// ===== 双三次插值子计算(__simt_callee__,拆分以满足50行限制)===== | ||
| 124 | + | ||
| 125 | +template <typename T> | ||
| 126 | +__simt_callee__ __aicore__ inline float ComputeBicubicPixel( | ||
| 127 | + __gm__ T* input, int64_t ncBase, | ||
| 128 | + int32_t inputHeight, int32_t inputWidth, | ||
| 129 | + int32_t inY, int32_t inX, | ||
| 130 | + float tX, float tY) | ||
| 131 | +{ | ||
| 132 | + float coefficients[BICUBIC_WINDOW_SIZE]; | ||
| 133 | + for (int32_t k = 0; k < BICUBIC_WINDOW_SIZE; k++) { | ||
| 134 | + int32_t iy = inY - 1 + k; | ||
| 135 | + float v0 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX - 1); | ||
| 136 | + float v1 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 0); | ||
| 137 | + float v2 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 1); | ||
| 138 | + float v3 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 2); | ||
| 139 | + coefficients[k] = CubicInterp1d(v0, v1, v2, v3, tX); | ||
| 140 | + } | ||
| 141 | + return CubicInterp1d(coefficients[0], coefficients[1], | ||
| 142 | + coefficients[2], coefficients[3], tY); | ||
| 143 | +} | ||
| 144 | + | ||
| 145 | +// ===== SIMT VF Kernel(R001/R002/R003/R005 全部应用)===== | ||
| 146 | +// | ||
| 147 | +// 模板参数: | ||
| 148 | +// T - 数据类型(float / half / bfloat16_t) | ||
| 149 | +// AlignCorners - 编译期常量,消除循环内 if-else 分支(R001/R007) | ||
| 150 | +// UintT - 索引/快除统一类型(uint32_t 或 uint64_t)(R003) | ||
| 151 | +// | ||
| 152 | +// 坐标拆解策略(R005 + R002): | ||
| 153 | +// N*C 合轴为 nc → 3D 拆解 (nc, oy, ox) | ||
| 154 | +// 仅 2 次 Simt::UintDiv(outputWidth, outputHeight),比原始 4D 少 1 次 | ||
| 155 | +// ncBase = nc * inputHeight * inputWidth,无需再分解 n 和 c | ||
| 156 | + | ||
| 157 | +template <typename T, bool AlignCorners, typename UintT> | ||
| 158 | +__simt_vf__ __aicore__ __launch_bounds__(BICUBIC_THREADS<UintT>) | ||
| 159 | +inline void OpUpsampleBicubic2dSimtKernel( | ||
| 160 | + UintT totalElements, | ||
| 161 | + int32_t ncSize, | ||
| 162 | + int32_t inputHeight, | ||
| 163 | + int32_t inputWidth, | ||
| 164 | + int32_t outputHeight, | ||
| 165 | + int32_t outputWidth, | ||
| 166 | + float scaleH, | ||
| 167 | + float scaleW, | ||
| 168 | + UintT owMagic, | ||
| 169 | + UintT owShift, | ||
| 170 | + UintT ohMagic, | ||
| 171 | + UintT ohShift, | ||
| 172 | + __gm__ T* input, | ||
| 173 | + __gm__ T* output) | ||
| 174 | +{ | ||
| 175 | + for (UintT idx = static_cast<UintT>(blockIdx.x) * static_cast<UintT>(blockDim.x) + static_cast<UintT>(threadIdx.x); | ||
| 176 | + idx < totalElements; | ||
| 177 | + idx += static_cast<UintT>(blockDim.x) * static_cast<UintT>(gridDim.x)) { | ||
| 178 | + | ||
| 179 | + // 3D 坐标拆解 (nc, oy, ox) — R005 合轴 + R002 快除 | ||
| 180 | + // idx = nc * (outputHeight * outputWidth) + oy * outputWidth + ox | ||
| 181 | + UintT ox = idx - Simt::UintDiv<UintT>(idx, owMagic, owShift) * static_cast<UintT>(outputWidth); | ||
| 182 | + UintT tmp = Simt::UintDiv<UintT>(idx, owMagic, owShift); | ||
| 183 | + UintT oy = tmp - Simt::UintDiv<UintT>(tmp, ohMagic, ohShift) * static_cast<UintT>(outputHeight); | ||
| 184 | + UintT nc = Simt::UintDiv<UintT>(tmp, ohMagic, ohShift); | ||
| 185 | + | ||
| 186 | + // 短路优化:输入输出尺寸相同时直接拷贝 | ||
| 187 | + if (inputHeight == outputHeight && inputWidth == outputWidth) { | ||
| 188 | + int64_t inIdx = static_cast<int64_t>(nc) * inputHeight * inputWidth + | ||
| 189 | + static_cast<int64_t>(oy) * inputWidth + static_cast<int64_t>(ox); | ||
| 190 | + output[idx] = input[inIdx]; | ||
| 191 | + continue; | ||
| 192 | + } | ||
| 193 | + | ||
| 194 | + // 计算源坐标 — AlignCorners 为编译期常量,if constexpr 零运行期开销(R001/R007) | ||
| 195 | + float realX, realY; | ||
| 196 | + if constexpr (AlignCorners) { | ||
| 197 | + realX = scaleW * static_cast<float>(ox); | ||
| 198 | + realY = scaleH * static_cast<float>(oy); | ||
| 199 | + } else { | ||
| 200 | + realX = scaleW * (static_cast<float>(ox) + HALF_PIXEL_OFFSET) - HALF_PIXEL_OFFSET; | ||
| 201 | + realY = scaleH * (static_cast<float>(oy) + HALF_PIXEL_OFFSET) - HALF_PIXEL_OFFSET; | ||
| 202 | + } | ||
| 203 | + | ||
| 204 | + int32_t inX = static_cast<int32_t>(floorf(realX)); | ||
| 205 | + float tX = realX - static_cast<float>(inX); | ||
| 206 | + int32_t inY = static_cast<int32_t>(floorf(realY)); | ||
| 207 | + float tY = realY - static_cast<float>(inY); | ||
| 208 | + | ||
| 209 | + // 双三次插值 — ncBase 直接用合轴后的 nc 计算,无需分解 n 和 c(R005) | ||
| 210 | + int64_t ncBase = static_cast<int64_t>(nc) * inputHeight * inputWidth; | ||
| 211 | + | ||
| 212 | + float result = ComputeBicubicPixel<T>( | ||
| 213 | + input, ncBase, inputHeight, inputWidth, | ||
| 214 | + inY, inX, tX, tY); | ||
| 215 | + | ||
| 216 | + output[idx] = FromFloat<T>(result); | ||
| 217 | + } | ||
| 218 | +} | ||
| 219 | + | ||
| 220 | +// ===== Process 入口(R003: 按数据量分派 32/64 位路径)===== | ||
| 221 | + | ||
| 222 | +template <typename T, bool AlignCorners> | ||
| 223 | +__aicore__ inline void Process(GM_ADDR x, GM_ADDR y, | ||
| 224 | + const UpsampleBicubic2dRegBaseTilingData* tilingData) | ||
| 225 | +{ | ||
| 226 | + __gm__ T* inputGm = (__gm__ T*)x; | ||
🟢 Low |
对比
upsample_bicubic2d 的 Process 缺少 实际影响:
修复建议:
![]() ![]() | |||
| 227 | + __gm__ T* outputGm = (__gm__ T*)y; | ||
| 228 | + | ||
| 229 | + if (tilingData->totalElements <= static_cast<int64_t>(INT32_MAX)) { | ||
| 230 | + // 32 位路径:索引运算吞吐翻倍(R003) | ||
| 231 | + uint32_t owMagic = 0, owShift = 0; | ||
| 232 | + GetUintDivMagicAndShift<uint32_t>(owMagic, owShift, static_cast<uint32_t>(tilingData->outputWidth)); | ||
| 233 | + uint32_t ohMagic = 0, ohShift = 0; | ||
| 234 | + GetUintDivMagicAndShift<uint32_t>(ohMagic, ohShift, static_cast<uint32_t>(tilingData->outputHeight)); | ||
| 235 | + | ||
| 236 | + asc_vf_call<OpUpsampleBicubic2dSimtKernel<T, AlignCorners, uint32_t>>( | ||
| 237 | + dim3(BICUBIC_THREADS<uint32_t>), | ||
| 238 | + static_cast<uint32_t>(tilingData->totalElements), | ||
| 239 | + tilingData->ncSize, | ||
| 240 | + tilingData->inputHeight, | ||
| 241 | + tilingData->inputWidth, | ||
| 242 | + tilingData->outputHeight, | ||
| 243 | + tilingData->outputWidth, | ||
| 244 | + tilingData->scaleH, | ||
| 245 | + tilingData->scaleW, | ||
| 246 | + owMagic, owShift, ohMagic, ohShift, | ||
| 247 | + inputGm, outputGm); | ||
| 248 | + } else { | ||
| 249 | + // 64 位路径:大数据量回退 | ||
| 250 | + uint64_t owMagic = 0, owShift = 0; | ||
| 251 | + GetUintDivMagicAndShift<uint64_t>(owMagic, owShift, static_cast<uint64_t>(tilingData->outputWidth)); | ||
| 252 | + uint64_t ohMagic = 0, ohShift = 0; | ||
| 253 | + GetUintDivMagicAndShift<uint64_t>(ohMagic, ohShift, static_cast<uint64_t>(tilingData->outputHeight)); | ||
| 254 | + | ||
| 255 | + asc_vf_call<OpUpsampleBicubic2dSimtKernel<T, AlignCorners, uint64_t>>( | ||
| 256 | + dim3(BICUBIC_THREADS<uint64_t>), | ||
| 257 | + static_cast<uint64_t>(tilingData->totalElements), | ||
| 258 | + static_cast<uint64_t>(tilingData->ncSize), | ||
| 259 | + static_cast<uint64_t>(tilingData->inputHeight), | ||
| 260 | + static_cast<uint64_t>(tilingData->inputWidth), | ||
| 261 | + static_cast<uint64_t>(tilingData->outputHeight), | ||
| 262 | + static_cast<uint64_t>(tilingData->outputWidth), | ||
| 263 | + tilingData->scaleH, | ||
| 264 | + tilingData->scaleW, | ||
| 265 | + owMagic, owShift, ohMagic, ohShift, | ||
| 266 | + inputGm, outputGm); | ||
| 267 | + } | ||
| 268 | +} | ||
| 269 | + | ||
| 270 | +} // namespace NsUpsampleBicubic2d | ||
| 271 | + | ||
| 272 | + | ||
| @@ -0,0 +1,33 @@ | |||
| 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 upsample_bicubic2d_tiling_data.h | ||
| 15 | + * \brief tiling data struct for upsample_bicubic2d | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +struct UpsampleBicubic2dRegBaseTilingData { | ||
| 22 | + int32_t needCoreNum; // 实际启动核数 | ||
| 23 | + int64_t totalElements; // 总输出元素数: N * C * H_out * W_out | ||
| 24 | + int32_t ncSize; // N * C(合轴后的批次-通道合并维度) | ||
| 25 | + int32_t inputHeight; // H_in | ||
| 26 | + int32_t inputWidth; // W_in | ||
| 27 | + int32_t outputHeight; // H_out | ||
| 28 | + int32_t outputWidth; // W_out | ||
| 29 | + float scaleH; // 高度方向缩放因子(已计算好的实际值) | ||
| 30 | + float scaleW; // 宽度方向缩放因子(已计算好的实际值) | ||
| 31 | +}; | ||
| 32 | + | ||
| 33 | + | ||
| @@ -0,0 +1,39 @@ | |||
| 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 upsample_bicubic2d_tiling_key.h | ||
| 15 | + * \brief tiling key declare for upsample_bicubic2d | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +// 编码 alignCorners 场景模式,dtype 由 DTYPE_ 宏自动实例化 | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +ASCENDC_TPL_ARGS_DECL( | ||
| 28 | + UpsampleBicubic2d, | ||
| 29 | + ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, | ||
| 30 | + UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE, | ||
| 31 | + UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE)); | ||
| 32 | + | ||
| 33 | +ASCENDC_TPL_SEL( | ||
| 34 | + ASCENDC_TPL_ARGS_SEL( | ||
| 35 | + ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, | ||
| 36 | + UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE, | ||
| 37 | + UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE))); | ||
| 38 | + | ||
| 39 | + | ||
| @@ -0,0 +1,38 @@ | |||
| 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 upsample_bicubic2d_apt.cpp | ||
| 15 | + * \brief arch35 SIMT kernel entry for upsample_bicubic2d | ||
| 16 | + */ | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +using namespace AscendC; | ||
| 25 | + | ||
| 26 | +template <uint32_t schMode> | ||
| 27 | +__global__ __aicore__ void upsample_bicubic2d(GM_ADDR x, GM_ADDR y, | ||
| 28 | + GM_ADDR workspace, GM_ADDR tiling) | ||
| 29 | +{ | ||
| 30 | + REGISTER_TILING_DEFAULT(UpsampleBicubic2dRegBaseTilingData); | ||
| 31 | + GET_TILING_DATA_WITH_STRUCT(UpsampleBicubic2dRegBaseTilingData, tilingData, tiling); | ||
| 32 | + | ||
| 33 | + if constexpr (schMode == UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE) { | ||
| 34 | + NsUpsampleBicubic2d::Process<DTYPE_INPUT, true>(x, y, &tilingData); | ||
| 35 | + } else { | ||
| 36 | + NsUpsampleBicubic2d::Process<DTYPE_INPUT, false>(x, y, &tilingData); | ||
🟢 Low | APT 中
这是 对比:
正向建议:可以把这个最佳实践(用 ![]() ![]() | |||
| 37 | + } | ||
| 38 | +} | ||
| @@ -10,7 +10,11 @@ | |||
| 10 | 10 | ||
| 11 | file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | 11 | file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) |
| 12 | if(UT_TEST_ALL OR OP_HOST_UT) | 12 | if(UT_TEST_ALL OR OP_HOST_UT) |
| 13 | - add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | 13 | + if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950") |
| 14 | + add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35) | ||
| 15 | + else() | ||
| 16 | + add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | ||
| 17 | + endif() | ||
| 14 | endif() | 18 | endif() |
| 15 | 19 | ||
| 16 | if(NOT (UT_TEST_ALL OR OP_API_UT)) | 20 | if(NOT (UT_TEST_ALL OR OP_API_UT)) |
| @@ -0,0 +1,87 @@ | |||
| 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 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +using namespace std; | ||
| 20 | +using namespace ge; | ||
| 21 | + | ||
| 22 | +class UpsampleBicubic2dTiling : public testing::Test { | ||
| 23 | +protected: | ||
| 24 | + static void SetUpTestCase() { std::cout << "UpsampleBicubic2dTiling SetUp" << std::endl; } | ||
| 25 | + | ||
| 26 | + static void TearDownTestCase() { std::cout << "UpsampleBicubic2dTiling TearDown" << std::endl; } | ||
| 27 | +}; | ||
| 28 | + | ||
| 29 | +std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend950"}}; | ||
| 30 | + | ||
| 31 | +TEST_F(UpsampleBicubic2dTiling, upsample_bicubic2d_float32_align_false) | ||
| 32 | +{ | ||
| 33 | + struct UpsampleBicubic2dCompileInfo { | ||
| 34 | + } compileInfo; | ||
| 35 | + gert::TilingContextPara tilingContextPara( | ||
| 36 | + "UpsampleBicubic2d", | ||
| 37 | + { | ||
| 38 | + {{{1, 1, 4, 4}, {1, 1, 4, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}, // input x | ||
| 39 | + }, | ||
| 40 | + { | ||
| 41 | + {{{1, 1, 8, 8}, {1, 1, 8, 8}}, ge::DT_FLOAT, ge::FORMAT_ND}, // output y | ||
| 42 | + }, | ||
| 43 | + { | ||
| 44 | + gert::TilingContextPara::OpAttr( | ||
| 45 | + "output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(std::vector<int64_t>{8, 8})), | ||
| 46 | + gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)), | ||
| 47 | + gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)), | ||
| 48 | + gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)), | ||
| 49 | + }, | ||
| 50 | + &compileInfo, "Ascend950", | ||
| 51 | + 64, // number of cores | ||
| 52 | + 262144, // ub size | ||
| 53 | + 4096); // tiling data max size | ||
| 54 | + uint64_t expectTilingKey = 0; // ALIGN_CORNERS_FALSE | ||
| 55 | + string expectTilingData = "1 64 17179869185 34359738372 4539628424389459976 1056964608 "; | ||
| 56 | + std::vector<size_t> expectWorkspaces = {16777216}; | ||
| 57 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | ||
| 58 | +} | ||
| 59 | + | ||
| 60 | +TEST_F(UpsampleBicubic2dTiling, upsample_bicubic2d_float16_align_true) | ||
| 61 | +{ | ||
| 62 | + struct UpsampleBicubic2dCompileInfo { | ||
| 63 | + } compileInfo; | ||
| 64 | + gert::TilingContextPara tilingContextPara( | ||
| 65 | + "UpsampleBicubic2d", | ||
| 66 | + { | ||
| 67 | + {{{2, 3, 4, 4}, {2, 3, 4, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // input x | ||
| 68 | + }, | ||
| 69 | + { | ||
| 70 | + {{{2, 3, 8, 8}, {2, 3, 8, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // output y | ||
| 71 | + }, | ||
| 72 | + { | ||
| 73 | + gert::TilingContextPara::OpAttr( | ||
| 74 | + "output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(std::vector<int64_t>{8, 8})), | ||
| 75 | + gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 76 | + gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)), | ||
| 77 | + gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)), | ||
| 78 | + }, | ||
| 79 | + &compileInfo, "Ascend950", | ||
| 80 | + 64, // number of cores | ||
| 81 | + 262144, // ub size | ||
| 82 | + 4096); // tiling data max size | ||
| 83 | + uint64_t expectTilingKey = 1; // ALIGN_CORNERS_TRUE | ||
| 84 | + string expectTilingData = "1 384 17179869190 34359738372 4529334482997608456 1054567863 "; | ||
| 85 | + std::vector<size_t> expectWorkspaces = {16777216}; | ||
| 86 | + ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | ||
| 87 | +} | ||
| @@ -16,6 +16,7 @@ if (UT_TEST_ALL OR OP_KERNEL_UT) | |||
| 16 | # 算子自己的tiling文件路径 | 16 | # 算子自己的tiling文件路径 |
| 17 | set(upsample_bicubic2d_tiling_files | 17 | set(upsample_bicubic2d_tiling_files |
| 18 | ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling.cpp | 18 | ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling.cpp |
| 19 | + ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling_arch35.cpp | ||
| 19 | # ${elewise_common_tiling_files} | 20 | # ${elewise_common_tiling_files} |
| 20 | ) | 21 | ) |
| 21 | # 使用AddOpTestCase | 22 | # 使用AddOpTestCase |
| @@ -230,7 +230,7 @@ | |||
| 230 | {"name":"ReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, | 230 | {"name":"ReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, |
| 231 | {"name":"LeakyReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, | 231 | {"name":"LeakyReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, |
| 232 | {"name":"UpsampleNearest", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "kirinx90", "kirin9030"], "auto_sync" : true}, | 232 | {"name":"UpsampleNearest", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "kirinx90", "kirin9030"], "auto_sync" : true}, |
| 233 | - {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}}, | 233 | + {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend310b": true, "kirinx90": true, "kirin9030": true, "ascend950": false}}, |
| 234 | {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}, | 234 | {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}, |
| 235 | {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}}, | 235 | {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}}, |
| 236 | {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}, | 236 | {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}, |
| @@ -587,6 +587,7 @@ | |||
| 587 | {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, | 587 | {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""}, |
| 588 | {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""}, | 588 | {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""}, |
| 589 | {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""}, | 589 | {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""}, |
| 590 | + {"name":"ExtractGlimpseV2", "compute_units": ["ascend950"], "auto_sync": false}, | ||
| 590 | {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""}, | 591 | {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""}, |
| 591 | {"name":"DIoUGrad", "compute_units": ["ascend950"], "auto_sync": false}, | 592 | {"name":"DIoUGrad", "compute_units": ["ascend950"], "auto_sync": false}, |
| 592 | {"name":"GIoUGrad", "compute_units": ["ascend950"], "auto_sync": false} | 593 | {"name":"GIoUGrad", "compute_units": ["ascend950"], "auto_sync": false} |


新增算子需同步刷新op_list.md