Star547
370
代码
代码
Issues26
Pull Requests51
Actions
项目讨论
Wiki
分析
Star547
370
已合并
blend_image;background_replace;mrgba算子迁移 #199
hang创建于 1月16日
blend_image;background_replace;mrgba算子迁移 #199
已合并
hang创建于 1月16日
从dailydeath/ops-cvmaster合入到cann/ops-cvmaster
共 66 个文件变更+3459-0
Aobjdetect/background_replace/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aobjdetect/background_replace/doc/aclnnBackgroundReplace.md+226-0
@@ -0,0 +1,226 @@
1+# aclnnBackgroundReplace
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+ 
17+- 算子功能：
18+将输入的新的背景图片与已有图片进行融合，通过掩码的方式将背景替换为新的背景。
19+ 
20+- 计算公式：
21+
22+ $$
23+ out = bkg * (1 - mask) + src * mask
24+ $$
25+ 
26+## 函数原型
27+ 
28+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnBackgroundReplaceGetWorkspaceSize”接口获取入参并根据流程计算所需workspace大小，再调用“aclnnBackgroundReplace”接口执行计算。
29+ 
30+* `aclnnStatus aclnnBackgroundReplaceGetWorkspaceSize(const aclTensor* bkg, const aclTensor* src, const aclTensor* mask, const aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)`
31+* `aclnnStatus aclnnBackgroundReplace(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`
32+ 
33+## aclnnBackgroundReplaceGetWorkspaceSize
34+ 
35+- **参数说明：**
36+ * bkg(aclTensor*, 计算输入):Device侧的aclTensor，数据类型支持UINT8、FLOAT16，shape支持HWC(C=1、3)。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
37+ * src(aclTensor*, 计算输入): Device侧的aclTensor，数据类型支持UINT8、FLOAT16，shape支持HWC(C=1、3)。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
38+ * mask(aclTensor*, 计算输入):Device侧的aclTensor，数据类型支持FLOAT16，shape支持HWC(C=1)。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
39+ * out(aclTensor*, 计算输出): Device侧的aclTensor，数据类型支持UINT8、FLOAT16，shape支持HWC(C=1、3)，数据类型和shape与输入背景图片bkg一致，只支持连续Tensor, [数据格式](../../../docs/zh/context/数据格式.md)支持ND。
40+ * workspaceSize(uint64_t \*, 出参): 返回需要在Device侧申请的workspace大小。
41+ * executor(aclOpExecutor \*\*, 出参): 返回op执行器，包含了算子计算流程。
42+ 
43+- **返回值：**
44+ 
45+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
46+ 
47+```
48+第一段接口完成入参校验，出现以下场景时报错：
49+161001 (ACLNN_ERR_PARAM_NULLPTR)：1. 传入的bkg、src、mask或out为空指针
50+161002 (ACLNN_ERR_PARAM_INVALID)：1. bkg、src、mask或out的数据类型不在支持的范围内。
51+ 2. bkg、src、mask或out的数据格式不在支持的范围内。
52+ 3. bkg、src或out的数据类型不一致。
53+ 4. bkg、src、mask或out的数据格式不一致。
54+ 
55+```
56+ 
57+## aclnnBackgroundReplace
58+ 
59+- **参数说明：**
60+ * workspace(void \*, 入参): 在Device侧申请的workspace内存地址。
61+ * workspaceSize(uint64_t, 入参): 在Device侧申请的workspace大小，由第一段接口aclnnBackgroundReplaceGetWorkspaceSize获取。
62+ * executor(aclOpExecutor \*, 入参): op执行器，包含了算子计算流程。
63+ * stream(aclrtStream, 入参): 指定执行任务的Stream。
64+ 
65+- **返回值：**
66+ 
67+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
68+ 
69+## 约束说明
70+- 确定性计算：
71+ - aclnnBackgroundReplace默认确定性实现
72+ 
73+## 调用示例
74+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
75+```Cpp
76+#include <iostream>
77+#include <vector>
78+#include "acl/acl.h"
79+#include "aclnnop/aclnn_background_replace.h"
80+ 
81+#define CHECK_RET(cond, return_expr) \
82+ do { \
83+ if (!(cond)) { \
84+ return_expr; \
85+ } \
86+ } while (0)
87+ 
88+#define LOG_PRINT(message, ...) \
89+ do { \
90+ printf(message, ##__VA_ARGS__); \
91+ } while (0)
92+ 
93+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
94+ int64_t shapeSize = 1;
95+ for (auto i : shape) {
96+ shapeSize *= i;
97+ }
98+ return shapeSize;
99+}
100+ 
101+int Init(int32_t deviceId, aclrtStream* stream) {
102+ // 固定写法，资源初始化
103+ auto ret = aclInit(nullptr);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
105+ ret = aclrtSetDevice(deviceId);
106+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
107+ ret = aclrtCreateStream(stream);
108+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
109+ return 0;
110+}
111+ 
112+template <typename T>
113+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
114+ aclDataType dataType, aclTensor** tensor) {
115+ auto size = GetShapeSize(shape) * sizeof(T);
116+ // 调用aclrtMalloc申请Device侧内存
117+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
118+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
119+ 
120+ // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上
121+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
123+ 
124+ // 计算连续tensor的strides
125+ std::vector<int64_t> strides(shape.size(), 1);
126+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
127+ strides[i] = shape[i + 1] * strides[i + 1];
128+ }
129+ 
130+ // 调用aclCreateTensor接口创建aclTensor
131+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
132+ shape.data(), shape.size(), *deviceAddr);
133+ return 0;
134+}
135+ 
136+int main() {
137+ // 1. （固定写法）device/stream初始化，参考acl API手册
138+ // 根据自己的实际device填写deviceId
139+ int32_t deviceId = 0;
140+ aclrtStream stream;
141+ auto ret = Init(deviceId, &stream);
142+ // check根据自己的需要处理
143+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
144+ 
145+ // 2. 构造输入与输出，需要根据API的接口自定义构造
146+ std::vector<int64_t> bkgShape = {4, 2};
147+ std::vector<int64_t> srcShape = {4, 2};
148+ std::vector<int64_t> maskShape = {4, 2};
149+ std::vector<int64_t> outShape = {4, 2};
150+ void* bkgDeviceAddr = nullptr;
151+ void* srcDeviceAddr = nullptr;
152+ void* maskDeviceAddr = nullptr;
153+ void* outDeviceAddr = nullptr;
154+ aclTensor* bkg = nullptr;
155+ aclTensor* src = nullptr;
156+ aclTensor* mask = nullptr;
157+ aclTensor* out = nullptr;
158+ std::vector<uint8_t> bkgHostData = {0, 1, 2, 3, 4, 5, 6, 7, 8};
159+ std::vector<uint8_t> srcHostData = {1, 2, 3, 4, 5, 6, 7, 8, 9};
160+ std::vector<float> maskHostData = {1, 1, 1, 1, 1, 1, 1, 1};
161+ std::vector<uint8_t> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
162+ // 创建bkg aclTensor
163+ ret = CreateAclTensor(bkgHostData, bkgShape, &bkgDeviceAddr, aclDataType::ACL_UINT8, &bkg);
164+ CHECK_RET(ret == ACL_SUCCESS, return ret);
165+ // 创建src aclTensor
166+ ret = CreateAclTensor(srcHostData, srcShape, &srcDeviceAddr, aclDataType::ACL_UINT8, &src);
167+ CHECK_RET(ret == ACL_SUCCESS, return ret);
168+ // 创建mask aclTensor
169+ ret = CreateAclTensor(maskHostData, maskShape, &maskDeviceAddr, aclDataType::ACL_FLOAT16, &mask);
170+ CHECK_RET(ret == ACL_SUCCESS, return ret);
171+ // 创建out aclTensor
172+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_UINT8, &out);
173+ CHECK_RET(ret == ACL_SUCCESS, return ret);
174+ 
175+ uint64_t workspaceSize = 0;
176+ aclOpExecutor* executor;
177+ 
178+ // aclnnBackgroundReplace接口调用示例
179+ // 3. 调用CANN算子库API, 需要修改为具体的API名称
180+ // 调用aclnnBackgroundReplace第一段接口
181+ ret = aclnnBackgroundReplaceGetWorkspaceSize(bkg, src, mask, out, &workspaceSize, &executor);
182+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnBackgroundReplaceGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
183+ // 根据第一段接口计算出的workspaceSize申请device内存
184+ void* workspaceAddr = nullptr;
185+ if (workspaceSize > 0) {
186+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
187+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
188+ }
189+ // 调用aclnnBackgroundReplace第二段接口
190+ ret = aclnnBackgroundReplace(workspaceAddr, workspaceSize, executor, stream);
191+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnBackgroundReplace failed. ERROR: %d\n", ret); return ret);
192+ 
193+ // 4. （固定写法）同步等待任务执行结束
194+ ret = aclrtSynchronizeStream(stream);
195+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
196+ 
197+ // 5. 获取输出的值，将Device侧内存上的结果拷贝至Host侧，需要根据具体API的接口定义修改
198+ auto size = GetShapeSize(outShape);
199+ std::vector<uint8_t> resultData(size, 0);
200+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
201+ size * sizeof(uint8_t), ACL_MEMCPY_DEVICE_TO_HOST);
202+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
203+ for (int64_t i = 0; i < size; i++) {
204+ LOG_PRINT("result[%ld] is: %u\n", i, resultData[i]);
205+ }
206+ 
207+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
208+ aclDestroyTensor(bkg);
209+ aclDestroyTensor(src);
210+ aclDestroyTensor(mask);
211+ aclDestroyTensor(out);
212+ 
213+ // 7. 释放Device资源，需要根据具体API的接口定义修改
214+ aclrtFree(bkgDeviceAddr);
215+ aclrtFree(srcDeviceAddr);
216+ aclrtFree(maskDeviceAddr);
217+ aclrtFree(outDeviceAddr);
218+ if (workspaceSize > 0) {
219+ aclrtFree(workspaceAddr);
220+ }
221+ aclrtDestroyStream(stream);
222+ aclrtResetDevice(deviceId);
223+ aclFinalize();
224+ return 0;
225+}
226+```
Aobjdetect/background_replace/op_graph/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_op_graph_sources()
Aobjdetect/background_replace/op_graph/background_replace_proto.h+38-0
@@ -0,0 +1,38 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file mgrba_custom_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
16+#define OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+namespace ge
20+{
21+ /**
22+ * @brief Give transparency to the image.
23+ *
24+ * @par Inputs:
25+ * @li rgb: A tensor of the type DT_UINT8.
26+ * @li alpha:A tensor of the type DT_UINT8.
27+ *
28+ * @par Outputs:
29+ * @li dst: A tensor of the type DT_UINT8.
30+ */
31+ REG_OP(BackgroundReplace)
32+ .INPUT(bkg, TensorType({ DT_UINT8, DT_FLOAT16 }))
33+ .INPUT(src, TensorType({ DT_UINT8, DT_FLOAT16 }))
34+ .INPUT(mask, TensorType({ DT_FLOAT16, DT_FLOAT16 }))
35+ .OUTPUT(out, TensorType({ DT_UINT8, DT_FLOAT16 }))
36+ .OP_END_FACTORY_REG(BackgroundReplace)
37+} // namespace ge
38+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
Aobjdetect/background_replace/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_modules_sources(OPTYPE background_replace ACLNNTYPE aclnn)
Aobjdetect/background_replace/op_host/background_replace_def.cpp+45-0
@@ -0,0 +1,45 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file background_replace.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class BackgroundReplace : public OpDef {
19+ public:
20+ explicit BackgroundReplace(const char* name) : OpDef(name) {
21+ this->Input("bkg")
22+ .ParamType(REQUIRED)
23+ .DataType({ge::DT_FLOAT16, ge::DT_UINT8})
24+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
25+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
26+ this->Input("src")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_FLOAT16, ge::DT_UINT8})
29+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
31+ this->Input("mask")
32+ .ParamType(REQUIRED)
33+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16})
34+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
35+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
36+ this->Output("out")
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_FLOAT16, ge::DT_UINT8})
39+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
40+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
41+ this->AICore().AddConfig("ascend310p");
42+ }
43+};
44+OP_ADD(BackgroundReplace);
45+} // namespace ops
Aobjdetect/background_replace/op_host/background_replace_infershape.cpp+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file background_replace.cc
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+using namespace ge;
18+namespace ge {
19+static ge::graphStatus InferShape4BackgroundReplace(gert::InferShapeContext* context) {
20+ const gert::Shape* bkgShape = context->GetInputShape(0);
21+ gert::Shape* outShape = context->GetOutputShape(0);
22+ *outShape = *bkgShape;
23+ return ge::GRAPH_SUCCESS;
24+}
25+ 
26+IMPL_OP_INFERSHAPE(BackgroundReplace).InferShape(InferShape4BackgroundReplace);
27+} // namespace ops
Aobjdetect/background_replace/op_host/background_replace_tiling.cpp+65-0
@@ -0,0 +1,65 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file background_replace.cc
13+ * \brief BackgroundReplace算子Tiling入口
14+ */
15+#include "register/op_def_registry.h"
16+#include "register/op_impl_registry.h"
17+#include "platform/platform_info.h"
18+#include "log/log.h"
19+#include "tiling/tiling_api.h"
20+#include "util/math_util.h"
21+#include "background_replace_tiling.h"
22+ 
23+namespace optiling {
24+ 
25+static constexpr uint64_t TILING_KEY_HALF_C1 = 1;
26+static constexpr uint64_t TILING_KEY_UINT8_C1 = 2;
27+static constexpr uint64_t TILING_KEY_HALF_C3 = 3;
28+static constexpr uint64_t TILING_KEY_UINT8_C3 = 4;
29+ 
30+static constexpr uint32_t ASCEND_310P_BLOCK_DIM = 8;
31+ 
32+static ge::graphStatus TilingBackgroundReplace(gert::TilingContext* context) {
33+ TilingDataBackgroundReplace tiling;
34+ auto tensorBkg = context -> GetInputTensor(0);
35+ auto tensorMask = context -> GetInputTensor(2);
36+ uint32_t maskLength = tensorMask->GetShapeSize();
37+ uint32_t bkgLength = tensorBkg->GetShapeSize();
38+ auto bkgDataType = tensorBkg->GetDataType();
39+ uint64_t tiling_key = 0;
40+ if (maskLength == bkgLength && bkgDataType == ge::DT_FLOAT16) {
41+ tiling_key = TILING_KEY_HALF_C1;
42+ } else if (maskLength == bkgLength && bkgDataType == ge::DT_UINT8) {
43+ tiling_key = TILING_KEY_UINT8_C1;
44+ } else if(maskLength != bkgLength && bkgDataType == ge::DT_FLOAT16) {
45+ tiling_key = TILING_KEY_HALF_C3;
46+ } else if (maskLength != bkgLength && bkgDataType == ge::DT_UINT8) {
47+ tiling_key = TILING_KEY_UINT8_C3;
48+ }
49+ tiling.set_size(maskLength);
50+ context->SetTilingKey(tiling_key);
51+ context->SetBlockDim(ASCEND_310P_BLOCK_DIM);
52+ tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
53+ context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
54+ return ge::GRAPH_SUCCESS;
55+}
56+ 
57+ static ge::graphStatus TilingPrepareForBackgroundReplace(gert::TilingParseContext* context) {
58+ OP_LOGD("BackgroundReplace", "TilingPrepareForBackgroundReplace start.");
59+ return ge::GRAPH_SUCCESS;
60+ }
61+ struct BackgroundReplaceCompileInfo {};
62+ IMPL_OP_OPTILING(BackgroundReplace)
63+ .Tiling(TilingBackgroundReplace)
64+ .TilingParse<BackgroundReplaceCompileInfo>(TilingPrepareForBackgroundReplace);
65+}
Aobjdetect/background_replace/op_host/background_replace_tiling.h+28-0
@@ -0,0 +1,28 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file background_replace.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILT_IN_OP_TILING_RUNTIME_BACKGROUND_REPLACE_TILING_H
16+#define OPS_BUILT_IN_OP_TILING_RUNTIME_BACKGROUND_REPLACE_TILING_H
17+#include <cstdint>
18+#include "register/tilingdata_base.h"
19+#include "tiling/tiling_api.h"
20+ 
21+namespace optiling {
22+BEGIN_TILING_DATA_DEF(TilingDataBackgroundReplace)
23+ TILING_DATA_FIELD_DEF(uint32_t, size);
24+END_TILING_DATA_DEF;
25+ 
26+REGISTER_TILING_DATA_CLASS(BackgroundReplace, TilingDataBackgroundReplace)
27+}
28+#endif // OPS_BUILT_IN_OP_TILING_RUNTIME_BACKGROUND_REPLACE_TILING_H
Aobjdetect/background_replace/op_host/config/ascend310p/background_replace_binary.json+99-0
@@ -0,0 +1,99 @@
1+{
2+ "op_type": "BackgroundReplace",
3+ "op_list": [
4+ {
5+ "bin_filename": "BackgroundReplace_FP16",
6+ "inputs": [
7+ {
8+ "name": "bkg",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "src",
19+ "index": 1,
20+ "dtype": "float16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "mask",
29+ "index": 2,
30+ "dtype": "float16",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ]
36+ }
37+ ],
38+ "outputs": [
39+ {
40+ "name": "out",
41+ "index": 0,
42+ "dtype": "float16",
43+ "format": "ND",
44+ "paramType": "required",
45+ "shape": [
46+ -2
47+ ]
48+ }
49+ ]
50+ },
51+ {
52+ "bin_filename": "BackgroundReplace_UINT8",
53+ "inputs": [
54+ {
55+ "name": "bkg",
56+ "index": 0,
57+ "dtype": "uint8",
58+ "format": "ND",
59+ "paramType": "required",
60+ "shape": [
61+ -2
62+ ]
63+ },
64+ {
65+ "name": "src",
66+ "index": 1,
67+ "dtype": "uint8",
68+ "format": "ND",
69+ "paramType": "required",
70+ "shape": [
71+ -2
72+ ]
73+ },
74+ {
75+ "name": "mask",
76+ "index": 2,
77+ "dtype": "float16",
78+ "format": "ND",
79+ "paramType": "required",
80+ "shape": [
81+ -2
82+ ]
83+ }
84+ ],
85+ "outputs": [
86+ {
87+ "name": "out",
88+ "index": 0,
89+ "dtype": "uint8",
90+ "format": "ND",
91+ "paramType": "required",
92+ "shape": [
93+ -2
94+ ]
95+ }
96+ ]
97+ }
98+ ]
99+}
Aobjdetect/background_replace/op_host/config/ascend310p/background_replace_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选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+[BackgroundReplace]
13+default=0
Aobjdetect/background_replace/op_kernel/background_replace.cpp+254-0
@@ -0,0 +1,254 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file background_replace.cpp
13+ * \brief BackgroundReplace 算子 Kernel 入口.
14+ */
15+ 
16+#include "kernel_operator.h"
17+#include "vector_scheduler.h"
18+using namespace AscendC;
19+ 
20+constexpr int32_t BUFFER_NUM = 1;
21+constexpr int32_t CHANNEL_NUM = 3;
22+ 
23+template <typename T1, typename T2>
24+class KernelBackgroundReplaceC1 {
25+public:
26+ __aicore__ inline KernelBackgroundReplaceC1()
27+ {}
28+ __aicore__ inline void Init(GM_ADDR bkg, GM_ADDR src, GM_ADDR mask, GM_ADDR out, GM_ADDR workspace, size_t bufferNum,
29+ size_t bufferBytes, size_t gmIdx, size_t gmDataLen)
30+ {
31+ if (bufferBytes <= 0) {
32+ return;
33+ }
34+ pipe.InitBuffer(inQueueBkg, bufferNum, bufferBytes);
35+ pipe.InitBuffer(inQueueSrc, bufferNum, bufferBytes);
36+ pipe.InitBuffer(outQueuedst, bufferNum, bufferBytes);
37+ if (sizeof(T1) == 1) {
38+ pipe.InitBuffer(calcBufX1, bufferBytes * sizeof(half));
39+ pipe.InitBuffer(calcBufX2, bufferBytes * sizeof(half));
40+ pipe.InitBuffer(inQueueMask, bufferNum, bufferBytes * sizeof(half));
41+ } else {
42+ pipe.InitBuffer(inQueueMask, bufferNum, bufferBytes);
43+ }
44+ 
45+ x1Gm.SetGlobalBuffer((__gm__ T1*)bkg + gmIdx, gmDataLen);
46+ x2Gm.SetGlobalBuffer((__gm__ T1*)src + gmIdx, gmDataLen);
47+ x3Gm.SetGlobalBuffer((__gm__ T2*)mask + gmIdx, gmDataLen);
48+ zGm.SetGlobalBuffer((__gm__ T1*)out + gmIdx, gmDataLen);
49+ }
50+ __aicore__ void CalcForAlign32(uint32_t idx, size_t len)
51+ {
52+ if (len <= 0) {
53+ return;
54+ }
55+ // copyIn
56+ auto bkgLocal = inQueueBkg.AllocTensor<T1>();
57+ auto srcLocal = inQueueSrc.AllocTensor<T1>();
58+ auto maskLocal = inQueueMask.AllocTensor<T2>();
59+ auto dstLocal = outQueuedst.AllocTensor<T1>();
60+ DataCopy(bkgLocal, x1Gm[idx], len);
61+ DataCopy(srcLocal, x2Gm[idx], len);
62+ DataCopy(maskLocal, x3Gm[idx], len);
63+ inQueueBkg.EnQue(bkgLocal);
64+ inQueueSrc.EnQue(srcLocal);
65+ inQueueMask.EnQue(maskLocal);
66+ 
67+ //compute
68+ bkgLocal = inQueueBkg.DeQue<T1>();
69+ srcLocal = inQueueSrc.DeQue<T1>();
70+ maskLocal = inQueueMask.DeQue<T2>();
71+ 
72+ if constexpr(sizeof(T1) == 1) {
73+ LocalTensor<half> bkgTmpLocal = calcBufX1.Get<half>();
74+ LocalTensor<half> srcTmpLocal = calcBufX2.Get<half>();
75+ Cast(bkgTmpLocal, bkgLocal, RoundMode::CAST_NONE, len);
76+ Cast(srcTmpLocal, srcLocal, RoundMode::CAST_NONE, len);
77+ 
78+ Mul(srcTmpLocal, srcTmpLocal, maskLocal, len);
79+ Mul(maskLocal, bkgTmpLocal, maskLocal, len);
80+ Sub(bkgTmpLocal, bkgTmpLocal, maskLocal, len);
81+ Add(bkgTmpLocal, bkgTmpLocal, srcTmpLocal, len);
82+ 
83+ Cast(dstLocal, bkgTmpLocal, RoundMode::CAST_NONE, len);
84+ } else {
85+ Mul(srcLocal, srcLocal, maskLocal, len);
86+ Mul(maskLocal, bkgLocal, maskLocal, len);
87+ Sub(bkgLocal, bkgLocal, maskLocal, len);
88+ Add(dstLocal, bkgLocal, srcLocal, len);
89+ }
90+ 
91+ //CopyOut
92+ outQueuedst.EnQue(dstLocal);
93+ inQueueBkg.FreeTensor(bkgLocal);
94+ inQueueSrc.FreeTensor(srcLocal);
95+ inQueueMask.FreeTensor(maskLocal);
96+ dstLocal = outQueuedst.DeQue<T1>();
97+ DataCopy(zGm[idx], dstLocal, len);
98+ outQueuedst.FreeTensor(dstLocal);
99+ }
100+ 
101+private:
102+ TPipe pipe;
103+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueBkg;
104+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueSrc;
105+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueMask;
106+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueuedst;
107+ TBuf<TPosition::VECCALC> calcBufX1;
108+ TBuf<TPosition::VECCALC> calcBufX2;
109+ GlobalTensor<T1> x1Gm;
110+ GlobalTensor<T1> x2Gm;
111+ GlobalTensor<T2> x3Gm;
112+ GlobalTensor<T1> zGm;
113+};
114+ 
115+template <typename T1, typename T2>
116+class KernelBackgroundReplaceC3 {
117+public:
118+ __aicore__ inline KernelBackgroundReplaceC3()
119+ {}
120+ __aicore__ inline void Init(GM_ADDR bkg, GM_ADDR src, GM_ADDR mask, GM_ADDR out, GM_ADDR workspace, size_t bufferNum,
121+ size_t bufferBytes, size_t gmIdx, size_t gmDataLen)
122+ {
123+ if (bufferBytes <= 0) {
124+ return;
125+ }
126+ size_t bkgBufferBytes = bufferBytes * CHANNEL_NUM; // src
127+ pipe.InitBuffer(inQueueBkg, bufferNum, bkgBufferBytes);
128+ pipe.InitBuffer(inQueueSrc, bufferNum, bkgBufferBytes);
129+ pipe.InitBuffer(outQueuedst, bufferNum, bkgBufferBytes);
130+ if (sizeof(T1) == 1) {
131+ pipe.InitBuffer(calcBufX1, bkgBufferBytes * sizeof(half));
132+ pipe.InitBuffer(calcBufX2, bkgBufferBytes * sizeof(half));
133+ // broadcast mask
134+ pipe.InitBuffer(inQueueMask, bufferNum, bufferBytes * sizeof(half));
135+ pipe.InitBuffer(calcBufX3, bkgBufferBytes * sizeof(half));
136+ } else {
137+ // broadcast mask
138+ pipe.InitBuffer(inQueueMask, bufferNum, bufferBytes);
139+ pipe.InitBuffer(calcBufX3, bkgBufferBytes);
140+ }
141+ 
142+ x1Gm.SetGlobalBuffer((__gm__ T1*)bkg + gmIdx * CHANNEL_NUM, gmDataLen * CHANNEL_NUM);
143+ x2Gm.SetGlobalBuffer((__gm__ T1*)src + gmIdx * CHANNEL_NUM, gmDataLen * CHANNEL_NUM);
144+ x3Gm.SetGlobalBuffer((__gm__ T2*)mask + gmIdx, gmDataLen);
145+ zGm.SetGlobalBuffer((__gm__ T1*)out + gmIdx * CHANNEL_NUM, gmDataLen * CHANNEL_NUM);
146+ }
147+ 
148+ __aicore__ void CalcForAlign32(uint32_t idx, size_t len)
149+ {
150+ if (len <= 0) {
151+ return;
152+ }
153+ size_t srclen = len * CHANNEL_NUM;
154+ // copyIn
155+ auto bkgLocal = inQueueBkg.AllocTensor<T1>();
156+ auto srcLocal = inQueueSrc.AllocTensor<T1>();
157+ auto maskLocal = inQueueMask.AllocTensor<T2>();
158+ 
159+ DataCopy(bkgLocal, x1Gm[idx * CHANNEL_NUM], srclen);
160+ DataCopy(srcLocal, x2Gm[idx * CHANNEL_NUM], srclen);
161+ DataCopy(maskLocal, x3Gm[idx], len);
162+ inQueueBkg.EnQue(bkgLocal);
163+ inQueueSrc.EnQue(srcLocal);
164+ inQueueMask.EnQue(maskLocal);
165+ 
166+ //compute
167+ bkgLocal = inQueueBkg.DeQue<T1>();
168+ srcLocal = inQueueSrc.DeQue<T1>();
169+ maskLocal = inQueueMask.DeQue<T2>();
170+ const uint32_t dimNum = 2;
171+ const uint32_t dstShape[dimNum] = {static_cast<uint32_t>(len), CHANNEL_NUM};
172+ const uint32_t srcShape[dimNum] = {static_cast<uint32_t>(len), 1};
173+ LocalTensor<half> maskC3Local = calcBufX3.Get<half>();
174+ auto dstLocal = outQueuedst.AllocTensor<T1>();
175+ BroadCast<half, dimNum, 1>(maskC3Local, maskLocal, dstShape, srcShape);
176+ if constexpr(sizeof(T1) == 1) {
177+ LocalTensor<half> bkgTmpLocal = calcBufX1.Get<half>();
178+ LocalTensor<half> srcTmpLocal = calcBufX2.Get<half>();
179+ Cast(bkgTmpLocal, bkgLocal, RoundMode::CAST_NONE, srclen);
180+ Cast(srcTmpLocal, srcLocal, RoundMode::CAST_NONE, srclen);
181+ 
182+ Mul(srcTmpLocal, srcTmpLocal, maskC3Local, srclen);
183+ Mul(maskC3Local, bkgTmpLocal, maskC3Local, srclen);
184+ Sub(bkgTmpLocal, bkgTmpLocal, maskC3Local, srclen);
185+ Add(bkgTmpLocal, bkgTmpLocal, srcTmpLocal, srclen);
186+ 
187+ Cast(dstLocal, bkgTmpLocal, RoundMode::CAST_NONE, srclen);
188+ } else {
189+ Mul(srcLocal, srcLocal, maskC3Local, srclen);
190+ Mul(maskC3Local, bkgLocal, maskC3Local, srclen);
191+ Sub(bkgLocal, bkgLocal, maskC3Local, srclen);
192+ Add(dstLocal, bkgLocal, srcLocal, srclen);
193+ }
194+ 
195+ //CopyOut
196+
197+ outQueuedst.EnQue(dstLocal);
198+ inQueueBkg.FreeTensor(bkgLocal);
199+ inQueueSrc.FreeTensor(srcLocal);
200+ inQueueMask.FreeTensor(maskLocal);
201+ dstLocal = outQueuedst.DeQue<T1>();
202+ DataCopy(zGm[idx * CHANNEL_NUM], dstLocal, srclen);
203+ outQueuedst.FreeTensor(dstLocal);
204+ }
205+protected:
206+ TPipe pipe;
207+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueBkg;
208+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueSrc;
209+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueMask;
210+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueuedst;
211+ TBuf<TPosition::VECCALC> calcBufX1;
212+ TBuf<TPosition::VECCALC> calcBufX2;
213+ TBuf<TPosition::VECCALC> calcBufX3;
214+ GlobalTensor<T1> x1Gm;
215+ GlobalTensor<T1> x2Gm;
216+ GlobalTensor<T2> x3Gm;
217+ GlobalTensor<T1> zGm;
218+};
219+ 
220+template <typename T1, typename T2>
221+__aicore__ void run_op(GM_ADDR bkg, GM_ADDR src, GM_ADDR mask, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling, float ubVarNum, bool isRGB=false)
222+{
223+ GET_TILING_DATA(tilingData, tiling);
224+ if (!isRGB) {
225+ VectorScheduler sch(tilingData.size, GetBlockNum(), BUFFER_NUM, ubVarNum, sizeof(T1));
226+ size_t orgVecIdx = GetBlockIdx() * sch.dataLenPerCore;
227+ KernelBackgroundReplaceC1<T1, T2> op;
228+ op.Init(bkg, src, mask, out, workspace, sch.bufferNum, sch.dataBytesPerLoop, orgVecIdx, sch.dataLen);
229+ sch.run(&op, sch.dataLen);
230+ } else {
231+ VectorScheduler sch(tilingData.size, GetBlockNum(), BUFFER_NUM, ubVarNum, sizeof(T1));
232+ size_t orgVecIdx = GetBlockIdx() * sch.dataLenPerCore;
233+ KernelBackgroundReplaceC3<T1, T2> op;
234+ op.Init(bkg, src, mask, out, workspace, sch.bufferNum, sch.dataBytesPerLoop, orgVecIdx, sch.dataLen);
235+ sch.run(&op, sch.dataLen);
236+ }
237+}
238+ 
239+extern "C" __global__ __aicore__ void background_replace(GM_ADDR bkg, GM_ADDR src, GM_ADDR mask, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
240+{
241+ if (TILING_KEY_IS(1)) {
242+ constexpr float ubVarNum = 5;
243+ run_op<half, half>(bkg, src, mask, out, workspace, tiling, ubVarNum);
244+ } else if (TILING_KEY_IS(2)) {
245+ constexpr float ubVarNum = 12;
246+ run_op<uint8_t, half>(bkg, src, mask, out, workspace, tiling, ubVarNum);
247+ } else if (TILING_KEY_IS(3)) {
248+ constexpr float ubVarNum = 100;
249+ run_op<half, half>(bkg, src, mask, out, workspace, tiling, ubVarNum, true);
250+ } else if (TILING_KEY_IS(4)) {
251+ constexpr float ubVarNum = 100;
252+ run_op<uint8_t, half>(bkg, src, mask, out, workspace, tiling, ubVarNum, true);
253+ }
254+}
Aobjdetect/background_replace/op_kernel/vector_scheduler.h+135-0
@@ -0,0 +1,135 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/*!
13+ * \file vector_scheduler.h
14+ * \brief
15+ */
16+#ifndef ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
17+#define ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
18+ 
19+#include "kernel_operator.h"
20+using namespace AscendC;
21+ 
22+constexpr size_t UB_SIZE_BYTE = 248 * 1024;
23+constexpr size_t ALIGN_SIZE_BYTES = 32;
24+constexpr size_t BLOCK_SIZE_BYTES = 32;
25+ 
26+__aicore__ inline size_t UpAlignN(size_t n, size_t N)
27+{
28+ if (N == 0) {
29+ return 0;
30+ }
31+ return (n + N - 1) / N * N;
32+}
33+ 
34+__aicore__ inline size_t DownAlignN(size_t n, size_t N)
35+{
36+ if (N == 0) {
37+ return 0;
38+ }
39+ return n / N * N;
40+}
41+ 
42+__aicore__ inline size_t UpAlign32(size_t n)
43+{
44+ return UpAlignN(n, ALIGN_SIZE_BYTES);
45+}
46+ 
47+__aicore__ inline size_t DownAlign32(size_t n)
48+{
49+ return DownAlignN(n, ALIGN_SIZE_BYTES);
50+}
51+ 
52+class VectorComputer {
53+public:
54+ __aicore__ inline VectorComputer() {};
55+ 
56+ __aicore__ inline void CalcForAlign32(uint32_t idx, size_t len) {};
57+};
58+ 
59+class VectorScheduler {
60+public:
61+ __aicore__ inline VectorScheduler(size_t contentLen, size_t blockDim, size_t bufferNum, float ubVarCount,
62+ size_t sizeofT)
63+ : blockDim(blockDim), bufferNum(bufferNum), ubVarCount(ubVarCount), sizeofT(sizeofT)
64+ {
65+ auto blockIdx = GetBlockIdx();
66+ this->dataLenPer32B = BLOCK_SIZE_BYTES / this->sizeofT;
67+ // L1
68+ this->dataLenPerCore = contentLen / this->blockDim;
69+ if (this->dataLenPerCore < this->dataLenPer32B) {
70+ this->dataLenPerCore = blockIdx == 0 ? contentLen : 0;
71+ this->dataLenTailL1 = 0;
72+ } else {
73+ this->dataLenTailL1 = contentLen % this->blockDim;
74+ }
75+ 
76+ // L2
77+ int maxUbSizePerVar = UB_SIZE_BYTE / ubVarCount / this->bufferNum;
78+ this->dataBytesPerLoop = DownAlign32(maxUbSizePerVar);
79+ this->dataLenPerLoop = this->dataBytesPerLoop / this->sizeofT;
80+ 
81+ this->dataLen = this->dataLenPerCore;
82+ if (blockIdx == this->blockDim - 1) {
83+ this->dataLen += this->dataLenTailL1;
84+ }
85+ this->bufferBytesPerVar = this->dataLen > this->dataLenPerLoop ? this->dataBytesPerLoop : UpAlign32(
86+ this->dataLen * this->sizeofT);
87+ }
88+ 
89+ template<class Computer>
90+ __aicore__ inline void run(Computer *computer, size_t len)
91+ {
92+ if (len <= 0) {
93+ return;
94+ }
95+ size_t loops = len / this->dataLenPerLoop;
96+ size_t tailLen = len % this->dataLenPerLoop;
97+ size_t tailLenA32 = DownAlignN(tailLen, this->dataLenPer32B);
98+ size_t tailLenBackoff = tailLen - tailLenA32;
99+ 
100+ uint32_t idx = 0;
101+ for (size_t i = 0; i < loops; i++) {
102+ computer->CalcForAlign32(idx, this->dataLenPerLoop);
103+ idx = idx + this->dataLenPerLoop;
104+ }
105+ if (tailLenA32) {
106+ idx = loops * this->dataLenPerLoop;
107+ computer->CalcForAlign32(idx, tailLenA32);
108+ }
109+ if (tailLenBackoff > 0) {
110+ idx = len >= this->dataLenPer32B ? len - this->dataLenPer32B : 0;
111+ computer->CalcForAlign32(idx, this->dataLenPer32B);
112+ }
113+ }
114+ 
115+public:
116+ float ubVarCount;
117+ size_t blockDim;
118+ size_t bufferNum;
119+ 
120+ size_t sizeofT;
121+ 
122+ size_t dataLenPer32B;
123+ // L1
124+ size_t dataLen;
125+ size_t dataLenPerCore;
126+ size_t dataLenTailL1;
127+ size_t bufferBytesPerVar;
128+ // L2
129+ size_t dataLenPerLoop;
130+ size_t dataBytesPerLoop;
131+ size_t loopL2;
132+ size_t dataLenTailL2;
133+};
134+ 
135+#endif // ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
Aobjdetect/background_replace/tests/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT OR OP_KERNEL_UT))
13+ list(REMOVE_ITEM CURRENT_DIRS ut)
14+endif()
15+ 
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aobjdetect/background_replace/tests/ut/CMakeLists.txt+27-0
@@ -0,0 +1,27 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+ 
14+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT))
15+ list(REMOVE_ITEM CURRENT_DIRS op_host)
16+endif()
17+ 
18+if(NOT (UT_TEST_ALL OR OP_KERNEL_UT))
19+ list(REMOVE_ITEM CURRENT_DIRS op_kernel)
20+endif()
21+ 
22+ 
23+foreach(SUB_DIR ${CURRENT_DIRS})
24+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
25+ add_subdirectory(${SUB_DIR})
26+ endif()
27+endforeach()
Aobjdetect/background_replace/tests/ut/op_graph/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/background_replace/tests/ut/op_graph/fusion_pass/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/background_replace/tests/ut/op_host/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/background_replace/tests/ut/op_kernel/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_library(background_replace SHARED
12+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel/background_replace.cpp
13+ ./test_background_replace.cpp
14+)
15+ 
16+target_compile_options(background_replace PRIVATE -g -include ${CMAKE_CURRENT_SOURCE_DIR}/test_background_replace.h)
17+ 
18+target_link_libraries(background_replace PRIVATE
19+ $<BUILD_INTERFACE:intf_llt_pub_asan>
20+ -Wl,--whole-archive
21+ -Wl,--no-as-needed
22+ ${PRIVATE_L}
23+ -Wl,--as-needed
24+ -Wl,--no-whole-archive
25+)
Aobjdetect/background_replace/tests/ut/op_kernel/background_replace_data/gen_data.py+39-0
@@ -0,0 +1,39 @@
1+#!/usr/bin/python3
2+# -*- coding:utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+import os
13+import sys
14+import numpy as np
15+ 
16+ 
17+def gen_golden_data_simple(dtype, channels):
18+ if(channels == '1'):
19+ bkg = np.random.uniform(0, 15, (480, 640, 1)).astype(dtype)
20+ src = np.random.uniform(0, 15, (480, 640, 1)).astype(dtype)
21+ mask = np.random.uniform(0, 15, (480, 640, 1)).astype(np.float16)
22+ out = (bkg - bkg * mask + src * mask).astype(dtype)
23+ else:
24+ bkg = np.random.uniform(0, 15, (480, 640, 3)).astype(dtype)
25+ src = np.random.uniform(0, 15, (480, 640, 3)).astype(dtype)
26+ mask = np.random.uniform(0, 15, (480, 640, 1)).astype(np.float16)
27+ out = (bkg - bkg * mask + src * mask).astype(dtype)
28+ bkg.tofile("./bkg.bin")
29+ src.tofile("./src.bin")
30+ mask.tofile("./mask.bin")
31+ out.tofile("./golden.bin")
32+ 
33+if __name__ == "__main__":
34+ datatype = sys.argv[1]
35+ if datatype == "uint8":
36+ dtype = np.uint8
37+ else:
38+ dtype = np.float16
39+ gen_golden_data_simple(dtype, sys.argv[2])
Aobjdetect/background_replace/tests/ut/op_kernel/test_background_replace.cpp+218-0
@@ -0,0 +1,218 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#include <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+#include "test_background_replace.h"
15+ 
16+#ifdef __CCE_KT_TEST__
17+#include "tikicpulib.h"
18+#include "../data_utils.h"
19+#include "string.h"
20+#include <iostream>
21+#include <string>
22+#endif
23+ 
24+#include <cstdint>
25+ 
26+using namespace std;
27+ 
28+ 
29+extern "C" __global__ __aicore__ void background_replace(GM_ADDR bkg, GM_ADDR src,GM_ADDR mask, GM_ADDR out,
30+ GM_ADDR workSpace, GM_ADDR tiling);
31+ 
32+class background_replace_test : public testing::Test {
33+ protected:
34+ static void SetUpTestCase() {
35+ cout << "background_replace_test SetUp\n" << endl;
36+ }
37+ static void TearDownTestCase() {
38+ cout << "background_replace_test TearDown\n" << endl;
39+ }
40+};
41+ 
42+TEST_F(background_replace_test, test_case_float16_c1) {
43+ size_t bkg_size = 480*640*2;
44+ size_t src_size = 480*640*2;
45+ size_t mask_size = 480*640*2;
46+ size_t out_size = 480*640*2;
47+ // inputs
48+ size_t tiling_data_size = sizeof(BackgroundReplaceTilingData);
49+ 
50+ uint8_t *bkg = (uint8_t*)AscendC::GmAlloc(bkg_size);
51+ uint8_t *src = (uint8_t*)AscendC::GmAlloc(src_size);
52+ uint8_t *mask = (uint8_t*)AscendC::GmAlloc(mask_size);
53+ uint8_t *golden = (uint8_t*)AscendC::GmAlloc(out_size);
54+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
55+ 
56+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
57+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
58+ uint32_t blockDim = 8;
59+ system("cp -r ../../../../../../../ops/objdetect/background_replace/tests/ut/op_kernel/background_replace_data ./");
60+ system("chmod -R 755 ./background_replace_data/");
61+ system("cd ./background_replace_data/ && rm -rf ./*bin");
62+ system("cd ./background_replace_data/ && python3 gen_data.py float16 1");
63+ char * path_ = get_current_dir_name();
64+ string path(path_);
65+ ReadFile(path + "/background_replace_data/bkg.bin", bkg_size, bkg, bkg_size);
66+ ReadFile(path + "/background_replace_data/src.bin", src_size, src, src_size);
67+ ReadFile(path + "/background_replace_data/mask.bin", mask_size, mask, mask_size);
68+ ReadFile(path + "/background_replace_data/golden.bin", out_size, golden, out_size);
69+ BackgroundReplaceTilingData* tilingDatafromBin = reinterpret_cast<BackgroundReplaceTilingData*>(tiling);
70+ tilingDatafromBin->size = 480*640;
71+ 
72+ ICPU_SET_TILING_KEY(1);
73+ ICPU_RUN_KF(background_replace, blockDim, bkg, src, mask, out, workspace, (uint8_t*)(tilingDatafromBin));
74+
75+ AscendC::GmFree(bkg);
76+ AscendC::GmFree(src);
77+ AscendC::GmFree(mask);
78+ AscendC::GmFree(out);
79+ AscendC::GmFree(golden);
80+ 
81+ AscendC::GmFree(workspace);
82+ AscendC::GmFree(tiling);
83+ free(path_);
84+}
85+ 
86+TEST_F(background_replace_test, test_case_u8_c1) {
87+ size_t bkg_size = 480*640;
88+ size_t src_size = 480*640;
89+ size_t mask_size = 480*640*2;
90+ size_t out_size = 480*640;
91+ // inputs
92+ size_t tiling_data_size = sizeof(BackgroundReplaceTilingData);
93+ 
94+ uint8_t *bkg = (uint8_t*)AscendC::GmAlloc(bkg_size);
95+ uint8_t *src = (uint8_t*)AscendC::GmAlloc(src_size);
96+ uint8_t *mask = (uint8_t*)AscendC::GmAlloc(mask_size);
97+ uint8_t *golden = (uint8_t*)AscendC::GmAlloc(out_size);
98+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
99+ 
100+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
101+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
102+ uint32_t blockDim = 8;
103+ system("cp -r ../../../../../../../ops/built-in/tests/ut/fast_op_test/background_replace/background_replace_data ./");
104+ system("chmod -R 755 ./background_replace_data/");
105+ system("cd ./background_replace_data/ && rm -rf ./*bin");
106+ system("cd ./background_replace_data/ && python3 gen_data.py uint8 1");
107+ char * path_ = get_current_dir_name();
108+ string path(path_);
109+ ReadFile(path + "/background_replace_data/bkg.bin", bkg_size, bkg, bkg_size);
110+ ReadFile(path + "/background_replace_data/src.bin", src_size, src, src_size);
111+ ReadFile(path + "/background_replace_data/mask.bin", mask_size, mask, mask_size);
112+ ReadFile(path + "/background_replace_data/golden.bin", out_size, golden, out_size);
113+ BackgroundReplaceTilingData* tilingDatafromBin = reinterpret_cast<BackgroundReplaceTilingData*>(tiling);
114+ tilingDatafromBin->size = 480*640;
115+ 
116+ ICPU_SET_TILING_KEY(2);
117+ ICPU_RUN_KF(background_replace, blockDim, bkg, src, mask, out, workspace, (uint8_t*)(tilingDatafromBin));
118+
119+ AscendC::GmFree(bkg);
120+ AscendC::GmFree(src);
121+ AscendC::GmFree(mask);
122+ AscendC::GmFree(out);
123+ AscendC::GmFree(golden);
124+ 
125+ AscendC::GmFree(workspace);
126+ AscendC::GmFree(tiling);
127+ free(path_);
128+}
129+ 
130+ 
131+ 
132+TEST_F(background_replace_test, test_case_float16_c3) {
133+ size_t bkg_size = 480*640*6;
134+ size_t src_size = 480*640*6;
135+ size_t mask_size = 480*640*2;
136+ size_t out_size = 480*640*6;
137+ // inputs
138+ size_t tiling_data_size = sizeof(BackgroundReplaceTilingData);
139+ 
140+ uint8_t *bkg = (uint8_t*)AscendC::GmAlloc(bkg_size);
141+ uint8_t *src = (uint8_t*)AscendC::GmAlloc(src_size);
142+ uint8_t *mask = (uint8_t*)AscendC::GmAlloc(mask_size);
143+ uint8_t *golden = (uint8_t*)AscendC::GmAlloc(out_size);
144+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
145+ 
146+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
147+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
148+ uint32_t blockDim = 8;
149+ system("cp -r ../../../../../../../ops/built-in/tests/ut/fast_op_test/background_replace/background_replace_data ./");
150+ system("chmod -R 755 ./background_replace_data/");
151+ system("cd ./background_replace_data/ && rm -rf ./*bin");
152+ system("cd ./background_replace_data/ && python3 gen_data.py float16 3");
153+ char * path_ = get_current_dir_name();
154+ string path(path_);
155+ ReadFile(path + "/background_replace_data/bkg.bin", bkg_size, bkg, bkg_size);
156+ ReadFile(path + "/background_replace_data/src.bin", src_size, src, src_size);
157+ ReadFile(path + "/background_replace_data/mask.bin", mask_size, mask, mask_size);
158+ ReadFile(path + "/background_replace_data/golden.bin", out_size, golden, out_size);
159+ BackgroundReplaceTilingData* tilingDatafromBin = reinterpret_cast<BackgroundReplaceTilingData*>(tiling);
160+ tilingDatafromBin->size = 480*640;
161+ 
162+ ICPU_SET_TILING_KEY(3);
163+ ICPU_RUN_KF(background_replace, blockDim, bkg, src, mask, out, workspace, (uint8_t*)(tilingDatafromBin));
164+
165+ AscendC::GmFree(bkg);
166+ AscendC::GmFree(src);
167+ AscendC::GmFree(mask);
168+ AscendC::GmFree(out);
169+ AscendC::GmFree(golden);
170+ 
171+ AscendC::GmFree(workspace);
172+ AscendC::GmFree(tiling);
173+ free(path_);
174+}
175+ 
176+TEST_F(background_replace_test, test_case_u8_c3) {
177+ size_t bkg_size = 480*640*3;
178+ size_t src_size = 480*640*3;
179+ size_t mask_size = 480*640*2;
180+ size_t out_size = 480*640*3;
181+ // inputs
182+ size_t tiling_data_size = sizeof(BackgroundReplaceTilingData);
183+ 
184+ uint8_t *bkg = (uint8_t*)AscendC::GmAlloc(bkg_size);
185+ uint8_t *src = (uint8_t*)AscendC::GmAlloc(src_size);
186+ uint8_t *mask = (uint8_t*)AscendC::GmAlloc(mask_size);
187+ uint8_t *golden = (uint8_t*)AscendC::GmAlloc(out_size);
188+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
189+ 
190+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
191+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
192+ uint32_t blockDim = 8;
193+ system("cp -r ../../../../../../../ops/built-in/tests/ut/fast_op_test/background_replace/background_replace_data ./");
194+ system("chmod -R 755 ./background_replace_data/");
195+ system("cd ./background_replace_data/ && rm -rf ./*bin");
196+ system("cd ./background_replace_data/ && python3 gen_data.py uint8 3");
197+ char * path_ = get_current_dir_name();
198+ string path(path_);
199+ ReadFile(path + "/background_replace_data/bkg.bin", bkg_size, bkg, bkg_size);
200+ ReadFile(path + "/background_replace_data/src.bin", src_size, src, src_size);
201+ ReadFile(path + "/background_replace_data/mask.bin", mask_size, mask, mask_size);
202+ ReadFile(path + "/background_replace_data/golden.bin", out_size, golden, out_size);
203+ BackgroundReplaceTilingData* tilingDatafromBin = reinterpret_cast<BackgroundReplaceTilingData*>(tiling);
204+ tilingDatafromBin->size = 480*640;
205+ 
206+ ICPU_SET_TILING_KEY(4);
207+ ICPU_RUN_KF(background_replace, blockDim, bkg, src, mask, out, workspace, (uint8_t*)(tilingDatafromBin));
208+
209+ AscendC::GmFree(bkg);
210+ AscendC::GmFree(src);
211+ AscendC::GmFree(mask);
212+ AscendC::GmFree(out);
213+ AscendC::GmFree(golden);
214+ 
215+ AscendC::GmFree(workspace);
216+ AscendC::GmFree(tiling);
217+ free(path_);
218+}
Aobjdetect/background_replace/tests/ut/op_kernel/test_background_replace.h+35-0
@@ -0,0 +1,35 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#ifndef _FAST_OP_TEST_BACKGROUND_REPLACE_TILING_H_
12+#define _FAST_OP_TEST_BACKGROUND_REPLACE_TILING_H_
13+ 
14+#include "kernel_tiling/kernel_tiling.h"
15+ 
16+struct BackgroundReplaceTilingData {
17+ uint32_t size;
18+};
19+ 
20+#pragma pack(1)
21+ 
22+#pragma pack()
23+ 
24+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
25+ __ubuf__ tilingStruct* tilingDataPointer = \
26+ reinterpret_cast<__ubuf__ tilingStruct*>((__ubuf__ uint8_t*)(tilingPointer));
27+ 
28+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
29+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
30+ 
31+#define GET_TILING_DATA(tilingData, tilingPointer) \
32+ BackgroundReplaceTilingData tilingData; \
33+ INIT_TILING_DATA(BackgroundReplaceTilingData, tilingDataPointer, tilingPointer); \
34+ (tilingData).size = tilingDataPointer->size;
35+#endif
Aobjdetect/blend_images_custom/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aobjdetect/blend_images_custom/doc/aclnnBlendImagesCustom.md+229-0
@@ -0,0 +1,229 @@
1+# aclnnBlendImagesCustom
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+ 
17+- 算子功能：完成张量rgb、frame和alpha的透明度乘法计算。
18+ 
19+- 计算公式：
20+ 
21+$$
22+out_{i*3}=rgb_{i*3} * (alpha_i / 255) + frame_{i*3}*(1 - alpha_i/255)
23+$$
24+ 
25+## 函数原型
26+ 
27+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnBlendImagesCustomGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnBlendImagesCustom”接口执行计算。
28+ 
29+* `aclnnStatus aclnnBlendImagesCustomGetWorkspaceSize(const aclTensor *rgb, const aclTensor *alpha, const aclTensor *frame, const aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)`
30+* `aclnnStatus aclnnBlendImagesCustom(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`
31+ 
32+## aclnnBlendImagesCustomGetWorkspaceSize
33+ 
34+- **参数说明：**
35+
36+ - rgb(aclTensor*, 计算输入): Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=3)，与alpha满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。只支持连续Tensor,[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
37+ - alpha(aclTensor*, 计算输入): Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=1)，与rgb满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。只支持连续Tensor,[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
38+ - frame(aclTensor*, 计算输入): Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=3)，与alpha满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。只支持连续Tensor,[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
39+ - out(aclTensor*, 计算输出): Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=3)，与frameshape一致。只支持连续Tensor,[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
40+ - workspaceSize(uint64_t\*, 出参): 返回需要在Device侧申请的workspace大小。
41+ - executor(aclOpExecutor\*\*, 出参): 返回op执行器，包含了算子计算流程。
42+
43+- **返回值：**
44+ 
45+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
46+ 
47+ ```
48+ 第一段接口完成入参校验，出现以下场景时报错：
49+ 返回161001 (ACLNN_ERR_PARAM_NULLPTR)：1. 传入的rgb、alpha、frame或out是空指针。
50+ 返回161002 (ACLNN_ERR_PARAM_INVALID)：1. rgb、alpha、frame的数据类型和数据格式不在支持的范围之内。
51+ 2. rgb、alpha、frame的shape无法做broadcast，rgb和frame支持HWC(C=3), alpha支持HWC(C=1)。
52+ ```
53+ 
54+## aclnnBlendImagesCustom
55+ 
56+- **参数说明：**
57+ - workspace(void \*, 入参): 在Device侧申请的workspace内存地址。
58+ - workspaceSize(uint64_t, 入参): 在Device侧申请的workspace大小，由第一段接口aclnnBlendImagesCustomGetWorkspaceSize获取。
59+ - executor(aclOpExecutor \*, 入参): op执行器，包含了算子计算流程。
60+ - stream(aclrtStream, 入参): 指定执行任务的Stream。
61+ 
62+- **返回值：**
63+ 
64+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
65+ 
66+## 约束说明
67+- 确定性计算：
68+ - aclnnBlendImagesCustom默认确定性实现
69+ 
70+## 调用示例
71+ 
72+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
73+ 
74+```Cpp
75+#include <iostream>
76+#include <vector>
77+#include "acl/acl.h"
78+#include "aclnnop/aclnn_blend_images_custom.h"
79+ 
80+#define CHECK_RET(cond, return_expr) \
81+ do { \
82+ if (!(cond)) { \
83+ return_expr; \
84+ } \
85+ } while (0)
86+ 
87+#define LOG_PRINT(message, ...) \
88+ do { \
89+ printf(message, ##__VA_ARGS__); \
90+ } while (0)
91+ 
92+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
93+ int64_t shapeSize = 1;
94+ for (auto i : shape) {
95+ shapeSize *= i;
96+ }
97+ return shapeSize;
98+}
99+ 
100+void PrintOutResult(std::vector<int64_t> &shape, void** deviceAddr) {
101+ auto size = GetShapeSize(shape);
102+ std::vector<float> resultData(size, 0);
103+ auto ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
104+ *deviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
105+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
106+ for (int64_t i = 0; i < size; i++) {
107+ LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]);
108+ }
109+}
110+ 
111+int Init(int32_t deviceId, aclrtStream* stream) {
112+ // 固定写法，资源初始化
113+ auto ret = aclInit(nullptr);
114+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
115+ ret = aclrtSetDevice(deviceId);
116+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
117+ ret = aclrtCreateStream(stream);
118+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
119+ return 0;
120+}
121+ 
122+template <typename T>
123+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
124+ aclDataType dataType, aclTensor** tensor) {
125+ auto size = GetShapeSize(shape) * sizeof(T);
126+ // 调用aclrtMalloc申请device侧内存
127+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
128+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
129+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
130+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
131+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
132+ 
133+ // 计算连续tensor的strides
134+ std::vector<int64_t> strides(shape.size(), 1);
135+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
136+ strides[i] = shape[i + 1] * strides[i + 1];
137+ }
138+ 
139+ // 调用aclCreateTensor接口创建aclTensor
140+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
141+ shape.data(), shape.size(), *deviceAddr);
142+ return 0;
143+}
144+ 
145+int main() {
146+ // 1. （固定写法）device/stream初始化，参考acl API手册
147+ // 根据自己的实际device填写deviceId
148+ int32_t deviceId = 0;
149+ aclrtStream stream;
150+ auto ret = Init(deviceId, &stream);
151+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
152+ 
153+ // 2. 构造输入与输出，需要根据API的接口自定义构造
154+ std::vector<int64_t> rgbShape = {4, 3};
155+ std::vector<int64_t> alphaShape = {4, 1};
156+ std::vector<int64_t> frameShape = {4, 3};
157+ std::vector<int64_t> outShape = {4, 3};
158+ 
159+ void* rgbDeviceAddr = nullptr;
160+ void* alphaDeviceAddr = nullptr;
161+ void* frameDeviceAddr = nullptr;
162+ void* outDeviceAddr = nullptr;
163+ 
164+ aclTensor* rgb = nullptr;
165+ aclTensor* alpha = nullptr;
166+ aclTensor* frame = nullptr;
167+ aclTensor* out = nullptr;
168+ 
169+ std::vector<float> rgbHostData = {10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120};
170+ std::vector<float> alphaHostData = {255, 255, 255, 255};
171+ std::vector<float> frameHostData = {10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120};
172+ std::vector<float> outHostData = {10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120};
173+ 
174+ ret = CreateAclTensor(rgbHostData, rgbShape, &rgbDeviceAddr, aclDataType::ACL_UINT8, &rgb);
175+ CHECK_RET(ret == ACL_SUCCESS, return ret);
176+ ret = CreateAclTensor(alphaHostData, alphaShape, &alphaDeviceAddr, aclDataType::ACL_UINT8, &alpha);
177+ CHECK_RET(ret == ACL_SUCCESS, return ret);
178+ ret = CreateAclTensor(frameHostData, frameShape, &frameDeviceAddr, aclDataType::ACL_UINT8, &frame);
179+ CHECK_RET(ret == ACL_SUCCESS, return ret);
180+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_UINT8, &out);
181+ CHECK_RET(ret == ACL_SUCCESS, return ret);
182+ 
183+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
184+ uint64_t workspaceSize = 0;
185+ aclOpExecutor* executor;
186+ 
187+ // 调用aclnnBlendImagesCustom第一段接口
188+ ret = aclnnBlendImagesCustomGetWorkspaceSize(rgb, alpha, frame, out, &workspaceSize, &executor);
189+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnBlendImagesCustomGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
190+ 
191+ // 根据第一段接口计算出的workspaceSize申请device内存
192+ void* workspaceAddr = nullptr;
193+ if (workspaceSize > 0) {
194+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
195+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
196+ }
197+ 
198+ // 调用aclnnBlendImagesCustom第二段接口
199+ ret = aclnnBlendImagesCustom(workspaceAddr, workspaceSize, executor, stream);
200+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnBlendImagesCustom failed. ERROR: %d\n", ret); return ret);
201+ 
202+ // 4. （固定写法）同步等待任务执行结束
203+ ret = aclrtSynchronizeStream(stream);
204+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
205+ 
206+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
207+ PrintOutResult(outShape, &outDeviceAddr);
208+ 
209+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
210+ aclDestroyTensor(rgb);
211+ aclDestroyTensor(alpha);
212+ aclDestroyTensor(frame);
213+ aclDestroyTensor(out);
214+ 
215+ // 7. 释放device资源
216+ aclrtFree(rgbDeviceAddr);
217+ aclrtFree(alphaDeviceAddr);
218+ aclrtFree(frameDeviceAddr);
219+ aclrtFree(outDeviceAddr);
220+ if (workspaceSize > 0) {
221+ aclrtFree(workspaceAddr);
222+ }
223+ aclrtDestroyStream(stream);
224+ aclrtResetDevice(deviceId);
225+ aclFinalize();
226+ 
227+ return 0;
228+}
229+```
Aobjdetect/blend_images_custom/op_graph/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_op_graph_sources()
Aobjdetect/blend_images_custom/op_graph/blend_images_custom_proto.h+39-0
@@ -0,0 +1,39 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_OTHER_H_
16+#define OPS_BUILT_IN_OP_PROTO_INC_NN_OTHER_H_
17+ 
18+#include "graph/operator_reg.h"
19+namespace ge
20+{
21+/**
22+* @brief Generate rgb and frame images into a out image with alpha transparency. \n
23+ 
24+* @par Inputs:
25+* @li rgb: A Int, dtype is uint8, rgb images data.
26+* @li alpha: A Int, dtype is uint8, alpha transparency images data.
27+* @li frame: A Int, dtype is uint8, frame images data. \n
28+ 
29+* @par Outputs:
30+* @li out: The out tensor. Dtype is same as rgb. \n
31+*/
32+REG_OP(BlendImagesCustom)
33+ .INPUT(rgb, TensorType({DT_UINT8}))
34+ .INPUT(alpha, TensorType({DT_UINT8}))
35+ .INPUT(frame, TensorType({DT_UINT8}))
36+ .OUTPUT(out, TensorType({DT_UINT8}))
37+ .OP_END_FACTORY_REG(BlendImagesCustom)
38+} // namespace ge
39+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_OTHER_H_
Aobjdetect/blend_images_custom/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_modules_sources(OPTYPE blend_images_custom ACLNNTYPE aclnn)
Aobjdetect/blend_images_custom/op_host/blend_images_custom_def.cpp+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class BlendImagesCustom : public OpDef {
19+public:
20+ explicit BlendImagesCustom(const char* name) : OpDef(name)
21+ {
22+ this->Input("rgb")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_UINT8})
25+ .Format({ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND});
27+ this->Input("alpha")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_UINT8})
30+ .Format({ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND});
32+ this->Input("frame")
33+ .ParamType(REQUIRED)
34+ .DataType({ge::DT_UINT8})
35+ .Format({ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND});
37+ this->Output("out")
38+ .ParamType(REQUIRED)
39+ .DataType({ge::DT_UINT8})
40+ .Format({ge::FORMAT_ND})
41+ .UnknownShapeFormat({ge::FORMAT_ND});
42+ 
43+ this->AICore().AddConfig("ascend310p");
44+ }
45+};
46+ 
47+OP_ADD(BlendImagesCustom);
48+}
Aobjdetect/blend_images_custom/op_host/blend_images_custom_infershape.cpp+34-0
@@ -0,0 +1,34 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom.cc
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+static const size_t FRAME_INDEX = 2;
19+static const size_t OUT_INDEX = 0;
20+ 
21+using namespace ge;
22+namespace ge {
23+static ge::graphStatus InferShape4BlendImagesCustom(gert::InferShapeContext *context) {
24+ // infer shape
25+ OP_LOGD(context->GetNodeName(), "Begin to do InferShape4BlendImagesCustom.");
26+ const gert::Shape *frame_shape = context->GetInputShape(FRAME_INDEX);
27+ gert::Shape *out_shape = context->GetOutputShape(OUT_INDEX);
28+ *out_shape = *frame_shape;
29+ OP_LOGD(context->GetNodeName(), "End to do InferShape4BlendImagesCustom.");
30+ return ge::GRAPH_SUCCESS;
31+}
32+ 
33+IMPL_OP_INFERSHAPE(BlendImagesCustom).InferShape(InferShape4BlendImagesCustom);
34+} // namespace ge
Aobjdetect/blend_images_custom/op_host/blend_images_custom_tiling.cpp+54-0
@@ -0,0 +1,54 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom_tiling.cpp
13+ * \brief BlendImagesCustom 算子 Tiling 入口.
14+ */
15+#include "blend_images_custom_tiling.h"
16+#include "register/op_def_registry.h"
17+#include "register/op_impl_registry.h"
18+#include "tiling/platform/platform_ascendc.h"
19+#include "util/math_util.h"
20+#include "log/log.h"
21+ 
22+using namespace optiling;
23+ 
24+static constexpr uint32_t ASCEND_310P_BLOCK_DIM = 8;
25+ 
26+static ge::graphStatus Tiling4BlendImagesCustom(gert::TilingContext* context) {
27+ if (context == nullptr) {
28+ return ge::GRAPH_FAILED;
29+ }
30+ OP_LOGD(context->GetNodeName(), "Tiling4BlendImagesCustom running begin");
31+ auto tensorAlpha = context->GetInputTensor(1);
32+ uint32_t totalAlphaLength = tensorAlpha->GetShapeSize();
33+ TilingDataBlendImages tiling_host;
34+ tiling_host.set_totalAlphaLength(totalAlphaLength);
35+ context->SetBlockDim(ASCEND_310P_BLOCK_DIM);
36+ tiling_host.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
37+ context->GetRawTilingData()->SetDataSize(tiling_host.GetDataSize());
38+ return ge::GRAPH_SUCCESS;
39+}
40+ 
41+static ge::graphStatus TilingPrepareForBlendImagesCustom(gert::TilingParseContext* context) {
42+ OP_LOGD("BlendImagesCustom", "TilingPrepareForBlendImagesCustom start.");
43+ if (context == nullptr) {
44+ return ge::GRAPH_SUCCESS;
45+ }
46+ return ge::GRAPH_SUCCESS;
47+}
48+struct BlendImagesCustomCompileInfo {};
49+ 
50+namespace optiling {
51+IMPL_OP_OPTILING(BlendImagesCustom)
52+ .Tiling(Tiling4BlendImagesCustom)
53+ .TilingParse<BlendImagesCustomCompileInfo>(TilingPrepareForBlendImagesCustom);
54+}
Aobjdetect/blend_images_custom/op_host/blend_images_custom_tiling.h+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom_tiling.h
13+ * \brief BlendImagesCustom 算子 TilingData 结构定义.
14+ */
15+ 
16+#ifndef OPS_BUILT_IN_OP_TILING_RUNTIME_BLEND_IMAGES_CUSTOM_TILING_H_
17+#define OPS_BUILT_IN_OP_TILING_RUNTIME_BLEND_IMAGES_CUSTOM_TILING_H_
18+ 
19+#include <cstdint>
20+#include "register/tilingdata_base.h"
21+#include "tiling/tiling_api.h"
22+ 
23+namespace optiling {
24+BEGIN_TILING_DATA_DEF(TilingDataBlendImages)
25+ TILING_DATA_FIELD_DEF(uint32_t, totalAlphaLength);
26+END_TILING_DATA_DEF;
27+ 
28+REGISTER_TILING_DATA_CLASS(BlendImagesCustom, TilingDataBlendImages)
29+}
30+ 
31+#endif // OPS_BUILT_IN_OP_TILING_RUNTIME_BLEND_IMAGES_CUSTOM_TILING_H_
Aobjdetect/blend_images_custom/op_host/config/ascend310p/blend_images_custom_binary.json+52-0
@@ -0,0 +1,52 @@
1+{
2+ "op_type": "BlendImagesCustom",
3+ "op_list": [
4+ {
5+ "bin_filename": "BlendImagesCustom_uint8",
6+ "inputs": [
7+ {
8+ "name": "rgb",
9+ "index": 0,
10+ "dtype": "uint8",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "alpha",
19+ "index": 1,
20+ "dtype": "uint8",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "frame",
29+ "index": 2,
30+ "dtype": "uint8",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ]
36+ }
37+ ],
38+ "outputs": [
39+ {
40+ "name": "out",
41+ "index": 0,
42+ "dtype": "uint8",
43+ "format": "ND",
44+ "paramType": "required",
45+ "shape": [
46+ -2
47+ ]
48+ }
49+ ]
50+ }
51+ ]
52+}
Aobjdetect/blend_images_custom/op_host/config/ascend310p/blend_images_custom_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选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+[BlendImagesCustom]
13+default=0
Aobjdetect/blend_images_custom/op_kernel/blend_images_custom.cpp+147-0
@@ -0,0 +1,147 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file blend_images_custom.cpp
13+ * \brief
14+ */
15+#include "kernel_tiling/kernel_tiling.h"
16+#include "kernel_operator.h"
17+#include "vector_scheduler.h"
18+using namespace AscendC;
19+ 
20+constexpr int32_t BUFFER_NUM = 1;
21+/* ratio: 1/255 = 0.003921568627451 */
22+constexpr float RATIO = 0.003921568627451;
23+constexpr int32_t LENGTH_RATIO = 3;
24+constexpr int32_t BROAD_CAST_DIM = 2;
25+constexpr float UB_VAR_NUM = 100;
26+ 
27+template <typename T>
28+class KernelBlendImages {
29+public:
30+ __aicore__ inline KernelBlendImages() {}
31+ __aicore__ inline void Init(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR frame, GM_ADDR out, size_t bufferNum, size_t bufferBytes,
32+ size_t gmIdx, size_t gmDataLen)
33+ {
34+ if (bufferBytes <= 0) {
35+ return;
36+ }
37+ pipe.InitBuffer(inQueueRgb, bufferNum, LENGTH_RATIO * bufferBytes);
38+ pipe.InitBuffer(inQueueAlpha, bufferNum, bufferBytes);
39+ pipe.InitBuffer(inQueueFrame, bufferNum, LENGTH_RATIO * bufferBytes);
40+ pipe.InitBuffer(outQueue, bufferNum, LENGTH_RATIO * bufferBytes);
41+ 
42+ pipe.InitBuffer(tmpBufferRgb, LENGTH_RATIO * bufferBytes * sizeof(half));
43+ pipe.InitBuffer(tmpBufferAlpha, bufferBytes * sizeof(half));
44+ pipe.InitBuffer(tmpBufferAlphaC3, LENGTH_RATIO * bufferBytes * sizeof(half));
45+ pipe.InitBuffer(tmpBufferFrame, LENGTH_RATIO * bufferBytes * sizeof(half));
46+ pipe.InitBuffer(tmpBufferFrameMulAlpha, LENGTH_RATIO * bufferBytes * sizeof(half));
47+ 
48+ rgbGm.SetGlobalBuffer((__gm__ T*)rgb + LENGTH_RATIO * gmIdx, LENGTH_RATIO * gmDataLen);
49+ alphaGm.SetGlobalBuffer((__gm__ T*)alpha + gmIdx, gmDataLen);
50+ frameGm.SetGlobalBuffer((__gm__ T*)frame + LENGTH_RATIO * gmIdx, LENGTH_RATIO * gmDataLen);
51+ outGm.SetGlobalBuffer((__gm__ T*)out + LENGTH_RATIO * gmIdx, LENGTH_RATIO * gmDataLen);
52+ }
53+ 
54+ __aicore__ inline void CalcForAlign32(uint32_t idx, size_t len)
55+ {
56+ uint32_t alphaIdx = idx;
57+ uint32_t rgbIdx = LENGTH_RATIO * idx;
58+ size_t alphaLen = len;
59+ size_t rgbLen = LENGTH_RATIO * len;
60+ if (len <= 0) {
61+ return ;
62+ }
63+ // copyIn
64+ auto rgbLocal = inQueueRgb.AllocTensor<T>();
65+ auto alphaLocal = inQueueAlpha.AllocTensor<T>();
66+ auto frameLocal = inQueueFrame.AllocTensor<T>();
67+ DataCopy(rgbLocal, rgbGm[rgbIdx], rgbLen);
68+ DataCopy(alphaLocal, alphaGm[alphaIdx], alphaLen);
69+ DataCopy(frameLocal, frameGm[rgbIdx], rgbLen);
70+ inQueueRgb.EnQue(rgbLocal);
71+ inQueueAlpha.EnQue(alphaLocal);
72+ inQueueFrame.EnQue(frameLocal);
73+ // compute
74+ rgbLocal = inQueueRgb.DeQue<T>();
75+ alphaLocal = inQueueAlpha.DeQue<T>();
76+ frameLocal = inQueueFrame.DeQue<T>();
77+ auto outLocal = outQueue.AllocTensor<T>();
78+ auto rgbHalfLocal = tmpBufferRgb.Get<half>();
79+ auto alphaHalfLocal = tmpBufferAlpha.Get<half>();
80+ auto alphaC3HalfLocal = tmpBufferAlphaC3.Get<half>();
81+ auto frameHalfLocal = tmpBufferFrame.Get<half>();
82+ auto frameMulAlphaHalfLocal = tmpBufferFrameMulAlpha.Get<half>();
83+ Cast(rgbHalfLocal, rgbLocal, RoundMode::CAST_NONE, rgbLen);
84+ Cast(alphaHalfLocal, alphaLocal, RoundMode::CAST_NONE, alphaLen);
85+ Cast(frameHalfLocal, frameLocal, RoundMode::CAST_NONE, rgbLen);
86+ half ratio = RATIO;
87+ Muls(alphaHalfLocal, alphaHalfLocal, ratio, alphaLen);
88+ const uint32_t dstShape[BROAD_CAST_DIM] = {static_cast<uint32_t>(alphaLen), LENGTH_RATIO};
89+ const uint32_t srcShape[BROAD_CAST_DIM] = {static_cast<uint32_t>(alphaLen), 1};
90+ BroadCast<half, BROAD_CAST_DIM, 1>(alphaC3HalfLocal, alphaHalfLocal, dstShape, srcShape);
91+ Mul(frameMulAlphaHalfLocal, frameHalfLocal, alphaC3HalfLocal, rgbLen);
92+ Sub(frameHalfLocal, frameHalfLocal, frameMulAlphaHalfLocal, rgbLen);
93+ Mul(rgbHalfLocal, rgbHalfLocal, alphaC3HalfLocal, rgbLen);
94+ Add(frameHalfLocal, frameHalfLocal, rgbHalfLocal, rgbLen);
95+ Cast(outLocal, frameHalfLocal, RoundMode::CAST_NONE, rgbLen);
96+ outQueue.EnQue<T>(outLocal);
97+ inQueueRgb.FreeTensor(rgbLocal);
98+ inQueueAlpha.FreeTensor(alphaLocal);
99+ inQueueFrame.FreeTensor(frameLocal);
100+ // CopyOut
101+ outLocal = outQueue.DeQue<T>();
102+ DataCopy(outGm[rgbIdx], outLocal, rgbLen);
103+ outQueue.FreeTensor(outLocal);
104+ }
105+ 
106+private:
107+ TPipe pipe;
108+ TBuf<QuePosition::VECCALC> tmpBufferRgb;
109+ TBuf<QuePosition::VECCALC> tmpBufferAlpha;
110+ TBuf<QuePosition::VECCALC> tmpBufferFrame;
111+ TBuf<QuePosition::VECCALC> tmpBufferAlphaC3;
112+ TBuf<QuePosition::VECCALC> tmpBufferFrameMulAlpha;
113+ // create queues for input, in this case depth is equal to buffer num
114+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueRgb;
115+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueAlpha;
116+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueFrame;
117+ // create queue for output, in this case depth is equal to buffer num
118+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueue;
119+ GlobalTensor<T> rgbGm;
120+ GlobalTensor<T> alphaGm;
121+ GlobalTensor<T> frameGm;
122+ GlobalTensor<T> outGm;
123+};
124+ 
125+template <typename T>
126+__aicore__ void run_op(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR frame, GM_ADDR out, GM_ADDR tiling, float ubVarNum) {
127+ GET_TILING_DATA(tilingData, tiling);
128+ VectorScheduler sch(tilingData.totalAlphaLength, GetBlockNum(), BUFFER_NUM, ubVarNum, sizeof(T));
129+ KernelBlendImages<T> op;
130+ size_t orgVecIdx = GetBlockIdx() * sch.dataLenPerCore;
131+ op.Init(rgb, alpha, frame, out, sch.bufferNum, sch.dataBytesPerLoop, orgVecIdx, sch.dataLen);
132+ sch.run(&op, sch.dataLen);
133+}
134+ 
135+extern "C" __global__ __aicore__ void blend_images_custom(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR frame, GM_ADDR out,
136+ GM_ADDR workspace, GM_ADDR tiling) {
137+ run_op<uint8_t>(rgb, alpha, frame, out, tiling, UB_VAR_NUM);
138+}
139+ 
140+#ifndef __CCE_KT_TEST__
141+// call of kernel function
142+void blend_images_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *rgb, uint8_t *alpha, uint8_t *frame,
143+ uint8_t *out, uint8_t *workspace, uint8_t *tiling)
144+{
145+ blend_images_custom<<<blockDim, l2ctrl, stream>>>(rgb, alpha, frame, out, workspace, tiling);
146+}
147+#endif
Aobjdetect/blend_images_custom/op_kernel/vector_scheduler.h+131-0
@@ -0,0 +1,131 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file vector_scheduler.h
13+ * \brief
14+ */
15+#ifndef ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
16+#define ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
17+ 
18+#include "kernel_operator.h"
19+using namespace AscendC;
20+ 
21+constexpr size_t UB_SIZE_BYTE = 248 * 1024;
22+constexpr size_t ALIGN_SIZE_BYTES = 32;
23+constexpr size_t BLOCK_SIZE_BYTES = 32;
24+ 
25+__aicore__ inline size_t UpAlignN(size_t n, size_t N)
26+{
27+ if (N == 0) {
28+ return 0;
29+ }
30+ return (n + N - 1) / N * N;
31+}
32+ 
33+__aicore__ inline size_t DownAlignN(size_t n, size_t N)
34+{
35+ if (N == 0) {
36+ return 0;
37+ }
38+ return n / N * N;
39+}
40+ 
41+__aicore__ inline size_t UpAlign32(size_t n)
42+{
43+ return UpAlignN(n, ALIGN_SIZE_BYTES);
44+}
45+ 
46+__aicore__ inline size_t DownAlign32(size_t n)
47+{
48+ return DownAlignN(n, ALIGN_SIZE_BYTES);
49+}
50+ 
51+class VectorComputer {
52+public:
53+ __aicore__ inline VectorComputer() {};
54+ 
55+ __aicore__ inline void CalcForAlign32(uint32_t idx, size_t len) {};
56+};
57+ 
58+class VectorScheduler {
59+public:
60+ __aicore__ inline VectorScheduler(size_t contentLen, size_t blockDim, size_t bufferNum, float ubVarCount,
61+ size_t sizeofT)
62+ : blockDim(blockDim), bufferNum(bufferNum), ubVarCount(ubVarCount), sizeofT(sizeofT)
63+ {
64+ auto blockIdx = GetBlockIdx();
65+ this->dataLenPer32B = BLOCK_SIZE_BYTES / this->sizeofT;
66+ // L1
67+ this->dataLenPerCore = contentLen / this->blockDim;
68+ if (this->dataLenPerCore < this->dataLenPer32B) {
69+ this->dataLenPerCore = blockIdx == 0 ? contentLen : 0;
70+ this->dataLenTailL1 = 0;
71+ } else {
72+ this->dataLenTailL1 = contentLen % this->blockDim;
73+ }
74+ // L2
75+ int maxUbSizePerVar = UB_SIZE_BYTE / ubVarCount / this->bufferNum;
76+ this->dataBytesPerLoop = DownAlign32(maxUbSizePerVar);
77+ this->dataLenPerLoop = this->dataBytesPerLoop / this->sizeofT;
78+ 
79+ this->dataLen = this->dataLenPerCore;
80+ if (blockIdx == this->blockDim - 1) {
81+ this->dataLen += this->dataLenTailL1;
82+ }
83+ this->bufferBytesPerVar = this->dataLen > this->dataLenPerLoop ? this->dataBytesPerLoop : UpAlign32(
84+ this->dataLen * this->sizeofT);
85+ }
86+ 
87+ template<class Computer>
88+ __aicore__ inline void run(Computer *computer, size_t len)
89+ {
90+ if (len <= 0) {
91+ return;
92+ }
93+ size_t loops = len / this->dataLenPerLoop;
94+ size_t tailLen = len % this->dataLenPerLoop;
95+ size_t tailLenA32 = DownAlignN(tailLen, this->dataLenPer32B);
96+ size_t tailLenBackoff = tailLen - tailLenA32;
97+ 
98+ uint32_t idx = 0;
99+ for (size_t i = 0; i < loops; i++) {
100+ computer->CalcForAlign32(idx, this->dataLenPerLoop);
101+ idx = idx + this->dataLenPerLoop;
102+ }
103+ if (tailLenA32) {
104+ idx = loops * this->dataLenPerLoop;
105+ computer->CalcForAlign32(idx, tailLenA32);
106+ }
107+ if (tailLenBackoff > 0) {
108+ idx = len >= this->dataLenPer32B ? len - this->dataLenPer32B : 0;
109+ computer->CalcForAlign32(idx, this->dataLenPer32B);
110+ }
111+ }
112+ 
113+public:
114+ float ubVarCount;
115+ size_t blockDim;
116+ size_t bufferNum;
117+ size_t sizeofT;
118+ size_t dataLenPer32B;
119+ // L1
120+ size_t dataLen;
121+ size_t dataLenPerCore;
122+ size_t dataLenTailL1;
123+ size_t bufferBytesPerVar;
124+ // L2
125+ size_t dataLenPerLoop;
126+ size_t dataBytesPerLoop;
127+ size_t loopL2;
128+ size_t dataLenTailL2;
129+};
130+ 
131+#endif // ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
Aobjdetect/blend_images_custom/tests/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT OR OP_KERNEL_UT))
13+ list(REMOVE_ITEM CURRENT_DIRS ut)
14+endif()
15+ 
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aobjdetect/blend_images_custom/tests/ut/CMakeLists.txt+27-0
@@ -0,0 +1,27 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+ 
14+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT))
15+ list(REMOVE_ITEM CURRENT_DIRS op_host)
16+endif()
17+ 
18+if(NOT (UT_TEST_ALL OR OP_KERNEL_UT))
19+ list(REMOVE_ITEM CURRENT_DIRS op_kernel)
20+endif()
21+ 
22+ 
23+foreach(SUB_DIR ${CURRENT_DIRS})
24+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
25+ add_subdirectory(${SUB_DIR})
26+ endif()
27+endforeach()
Aobjdetect/blend_images_custom/tests/ut/op_graph/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/blend_images_custom/tests/ut/op_graph/fusion_pass/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/blend_images_custom/tests/ut/op_host/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/blend_images_custom/tests/ut/op_kernel/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_library(blend_images_custom SHARED
12+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel/blend_images_custom.cpp
13+ ./test_blend_images_custom.cpp
14+)
15+ 
16+target_compile_options(blend_images_custom PRIVATE -g -include ${CMAKE_CURRENT_SOURCE_DIR}/test_blend_images_custom.h)
17+ 
18+target_link_libraries(blend_images_custom PRIVATE
19+ $<BUILD_INTERFACE:intf_llt_pub_asan>
20+ -Wl,--whole-archive
21+ -Wl,--no-as-needed
22+ ${PRIVATE_L}
23+ -Wl,--as-needed
24+ -Wl,--no-whole-archive
25+)
Aobjdetect/blend_images_custom/tests/ut/op_kernel/blend_images_custom/gen_data.py+27-0
@@ -0,0 +1,27 @@
1+#!/usr/bin/python3
2+# -*- coding:utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+import sys
13+import numpy as np
14+ 
15+def gen_golden_data_simple():
16+ rgb = np.random.uniform(0, 255, (480, 640, 3)).astype(np.uint8)
17+ alpha = np.random.uniform(0, 255, (480, 640, 1)).astype(np.uint8)
18+ frame = np.random.uniform(0, 255, (480, 640, 3)).astype(np.uint8)
19+ 
20+ out = np.uint8(np.float16(rgb) * np.float16(alpha) * 0.003921568627451 + np.float16(frame) - np.float16(frame) * np.float16(alpha) * 0.003921568627451)
21+ rgb.tofile("./rgb.bin")
22+ alpha.tofile("./alpha.bin")
23+ frame.tofile("./frame.bin")
24+ out.tofile("./out.bin")
25+ 
26+if __name__ == "__main__":
27+ gen_golden_data_simple()
Aobjdetect/blend_images_custom/tests/ut/op_kernel/test_blend_images_custom.cpp+114-0
@@ -0,0 +1,114 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#include <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+#include "test_blend_images_custom.h"
15+#ifdef __CCE_KT_TEST__
16+#include "tikicpulib.h"
17+#include "../data_utils.h"
18+#include "string.h"
19+#include <iostream>
20+#include <string>
21+#endif
22+ 
23+#include <cstdint>
24+ 
25+using namespace std;
26+ 
27+extern "C" __global__ __aicore__ void blend_images_custom(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR frame, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling);
28+class blend_images_custom_test : public testing::Test {
29+ protected:
30+ static void SetUpTestCase() {
31+ cout << "blend_images_custom_test SetUp\n" << endl;
32+ }
33+ static void TearDownTestCase() {
34+ cout << "blend_images_custom_test TearDown\n" << endl;
35+ }
36+};
37+ 
38+TEST_F(blend_images_custom_test, test_case_uint8) {
39+ size_t rgb_size = 480 * 640 * 3;
40+ size_t alpha_size = 480 * 640 * 1;
41+ size_t frame_size = 480 * 640 * 3;
42+ size_t out_size = 480 * 640 * 3;
43+ // inputs
44+ size_t tiling_data_size = sizeof(TilingDataBlendImages);
45+ 
46+ uint8_t *rgb = (uint8_t*)AscendC::GmAlloc(rgb_size);
47+ uint8_t *alpha = (uint8_t*)AscendC::GmAlloc(alpha_size);
48+ uint8_t *frame = (uint8_t*)AscendC::GmAlloc(frame_size);
49+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
50+
51+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
52+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
53+ uint32_t blockDim = 8;
54+ system("cp -r ../../../../../../../ops/objdetect/blend_images_custom/tests/ut/op_kernel/blend_images_custom_data ./");
55+ system("chmod -R 755 ./blend_images_custom_data/");
56+ system("cd ./blend_images_custom_data/ && rm -rf ./*bin");
57+ system("cd ./blend_images_custom_data/ && python3 gen_data.py");
58+ char * path_ = get_current_dir_name();
59+ string path(path_);
60+ ReadFile(path + "/blend_images_custom_data/rgb.bin", rgb_size, rgb, rgb_size);
61+ ReadFile(path + "/blend_images_custom_data/alpha.bin", alpha_size, alpha, alpha_size);
62+ ReadFile(path + "/blend_images_custom_data/frame.bin", frame_size, frame, frame_size);
63+ ReadFile(path + "/blend_images_custom_data/out.bin", out_size, out, out_size);
64+ TilingDataBlendImages* tilingDatafromBin = reinterpret_cast<TilingDataBlendImages*>(tiling);
65+ tilingDatafromBin->totalAlphaLength = 480 * 640 * 1;
66+
67+ ICPU_RUN_KF(blend_images_custom, blockDim, rgb, alpha, frame, out, workspace, (uint8_t*)(tilingDatafromBin));
68+
69+ AscendC::GmFree(rgb);
70+ AscendC::GmFree(alpha);
71+ AscendC::GmFree(frame);
72+ AscendC::GmFree(out);
73+ AscendC::GmFree(tiling);
74+ free(path_);
75+}
76+ 
77+TEST_F(blend_images_custom_test, test_case_uint8_size_is_zero) {
78+ size_t rgb_size = 0;
79+ size_t alpha_size = 0;
80+ size_t frame_size = 0;
81+ size_t out_size = 0;
82+ // inputs
83+ size_t tiling_data_size = sizeof(TilingDataBlendImages);
84+ 
85+ uint8_t *rgb = (uint8_t*)AscendC::GmAlloc(rgb_size);
86+ uint8_t *alpha = (uint8_t*)AscendC::GmAlloc(alpha_size);
87+ uint8_t *frame = (uint8_t*)AscendC::GmAlloc(frame_size);
88+ uint8_t *out = (uint8_t*)AscendC::GmAlloc(out_size);
89+
90+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
91+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
92+ uint32_t blockDim = 8;
93+ system("cp -r ../../../../../../../ops/built-in/tests/ut/fast_op_test/blend_images_custom/blend_images_custom_data ./");
94+ system("chmod -R 755 ./blend_images_custom_data/");
95+ system("cd ./blend_images_custom_data/ && rm -rf ./*bin");
96+ system("cd ./blend_images_custom_data/ && python3 gen_data.py");
97+ char * path_ = get_current_dir_name();
98+ string path(path_);
99+ ReadFile(path + "/blend_images_custom_data/rgb.bin", rgb_size, rgb, rgb_size);
100+ ReadFile(path + "/blend_images_custom_data/alpha.bin", alpha_size, alpha, alpha_size);
101+ ReadFile(path + "/blend_images_custom_data/frame.bin", frame_size, frame, frame_size);
102+ ReadFile(path + "/blend_images_custom_data/out.bin", out_size, out, out_size);
103+ TilingDataBlendImages* tilingDatafromBin = reinterpret_cast<TilingDataBlendImages*>(tiling);
104+ tilingDatafromBin->totalAlphaLength = 0;
105+
106+ ICPU_RUN_KF(blend_images_custom, blockDim, rgb, alpha, frame, out, workspace, (uint8_t*)(tilingDatafromBin));
107+
108+ AscendC::GmFree(rgb);
109+ AscendC::GmFree(alpha);
110+ AscendC::GmFree(frame);
111+ AscendC::GmFree(out);
112+ AscendC::GmFree(tiling);
113+ free(path_);
114+}
Aobjdetect/blend_images_custom/tests/ut/op_kernel/test_blend_images_custom.h+35-0
@@ -0,0 +1,35 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#ifndef _FAST_OP_TEST_BLEND_IMAGES_CUSTOM_TILING_H_
12+#define _FAST_OP_TEST_BLEND_IMAGES_CUSTOM_TILING_H_
13+ 
14+#include "register/tilingdata_base.h"
15+ 
16+struct TilingDataBlendImages {
17+ uint32_t totalAlphaLength = 0;
18+};
19+ 
20+#pragma pack(1)
21+ 
22+#pragma pack()
23+ 
24+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
25+ __ubuf__ tilingStruct* tilingDataPointer = \
26+ reinterpret_cast<__ubuf__ tilingStruct*>((__ubuf__ uint8_t*)(tilingPointer));
27+ 
28+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
29+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
30+ 
31+#define GET_TILING_DATA(tilingData, tilingPointer) \
32+ TilingDataBlendImages tilingData; \
33+ INIT_TILING_DATA(TilingDataBlendImages, tilingDataPointer, tilingPointer); \
34+ (tilingData).totalAlphaLength = tilingDataPointer->totalAlphaLength;
35+#endif // _FAST_OP_TEST_BLEND_IMAGES_CUSTOM_TILING_H_
Aobjdetect/mrgba_custom/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aobjdetect/mrgba_custom/doc/aclnnMrgbaCustom.md+216-0
@@ -0,0 +1,216 @@
1+# aclnnMrgbaCustom
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+- 算子功能：完成张量rgb和张量alpha的透明度乘法计算
17+ 
18+- 计算公式：out = rgb * ((broadcast)alpha/255)
19+ 
20+- 示例：
21+ 假设rgb是一张三通道彩色图片，alpha是其对应的透明度（单通道），使用该算子后可以将该图片生成带透明度的三通道图片。
22+ 
23+## 函数原型
24+ 
25+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用"aclnnMrgbaCustomGetWorkspaceSize"接口获取计算所需
26+workspace大小以及包含了算子计算流程的执行器，再调用"aclnnMrgbaCustom"接口执行计算。
27+ 
28+- `aclnnStatus aclnnMrgbaCustomGetWorkspaceSize(const aclTensor *rgb,const aclTensor *alpha, const aclTensor *out, uint64_t *workspaceSize,aclOpExecutor **executor)`
29+ 
30+- `aclnnStatus aclnnMrgbaCustom(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`
31+ 
32+## aclnnMrgbaCustomGetWorkspaceSize
33+ 
34+- **参数说明**：
35+ 
36+ - rgb(aclTensor*, 计算输入)：公式中的rgb，Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=3)，与alpha满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
37+ - alpha(aclTensor*, 计算输入)：公式中的alpha，Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=1)，与rgb满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
38+ - out(aclTensor*, 计算输出)：公式中的out，Device侧的aclTensor，数据类型支持UINT8，shape支持HWC(C=3)，与rgb的shape一致。只支持连续Tensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
39+ - workspaceSize(uint64_t*, 出参)：返回需要在Device侧申请的workspace大小。
40+ - executor(aclOpExecutor**, 出参)：返回op执行器，包含了算子计算流程。
41+- **返回值**：
42+ 
43+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
44+ 
45+ ```
46+ 第一段接口完成入参校验，出现以下场景时报错：
47+ 返回161001（ACLNN_ERR_PARAM_NULLPTR）：1.rgb、alpha或out是空指针。
48+ 返回161002（ACLNN_ERR_PARAM_INVALID）：1.rgb和alpha的数据类型不在支持的范围之内。
49+ 2.rgb和alpha的shape不满足HWC(C=3)和HWC(C=1)的要求。
50+ ```
51+ 
52+ 
53+## aclnnMrgbaCustom
54+ 
55+- **参数说明**：
56+ 
57+ - workspace(void*, 入参)：在Device侧申请的workspace内存地址。
58+ - workspaceSize(uint64_t, 入参)：在Device侧申请的workspace大小，由第一段接口aclnnMrgbaCustomGetWorkspaceSize获取。
59+ - executor(aclOpExecutor*, 入参)：op执行器，包含了算子计算流程。
60+ - stream(aclrtStream, 入参)：指定执行任务的Stream。
61+- **返回值**：
62+ 
63+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
64+ 
65+## 约束说明
66+- 确定性计算：
67+ - aclnnMrgbaCustom默认确定性实现
68+ 
69+## 调用示例
70+ 
71+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
72+ 
73+```Cpp
74+#include <iostream>
75+#include <vector>
76+#include "acl/acl.h"
77+#include "aclnnop/aclnn_mrgba_custom.h"
78+ 
79+#define CHECK_RET(cond, return_expr) \
80+ do { \
81+ if (!(cond)) { \
82+ return_expr; \
83+ } \
84+ } while (0)
85+ 
86+#define LOG_PRINT(message, ...) \
87+ do { \
88+ printf(message, ##__VA_ARGS__); \
89+ } while (0)
90+ 
91+int64_t GetShapeSize(const std::vector<int64_t> &shape)
92+{
93+ int64_t shapeSize = 1;
94+ for (auto i: shape) {
95+ shapeSize *= i;
96+ }
97+ return shapeSize;
98+}
99+ 
100+int Init(int32_t deviceId, aclrtStream *stream)
101+{
102+ // 固定写法，资源初始化
103+ auto ret = aclInit(nullptr);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
105+ ret = aclrtSetDevice(deviceId);
106+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
107+ ret = aclrtCreateStream(stream);
108+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
109+ return 0;
110+}
111+ 
112+template<typename T>
113+int CreateAclTensor(const std::vector <T> &hostData, const std::vector <int64_t> &shape, void **deviceAddr,
114+ aclDataType dataType, aclTensor **tensor)
115+{
116+ auto size = GetShapeSize(shape) * sizeof(T);
117+ // 调用aclrtMalloc申请device侧内存
118+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
120+ 
121+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
122+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
123+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
124+ 
125+ // 计算连续tensor的strides
126+ std::vector<int64_t> strides(shape.size(), 1);
127+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
128+ strides[i] = shape[i + 1] * strides[i + 1];
129+ }
130+ 
131+ // 调用aclCreateTensor接口创建aclTensor
132+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
133+ shape.data(), shape.size(), *deviceAddr);
134+ return 0;
135+}
136+ 
137+int main()
138+{
139+ // 1. (固定写法)device/stream初始化，参考acl对外接口列表
140+ // 根据自己的实际device填写deviceId
141+ int32_t deviceId = 0;
142+ aclrtStream stream;
143+ auto ret = Init(deviceId, &stream);
144+ // check根据自己的需要处理
145+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
146+ // 2.构造输入与输出，需要根据API的接口自定义构造
147+ std::vector<int64_t> rgbShape = {4, 3};
148+ std::vector<int64_t> alphaShape = {4, 1};
149+ std::vector<int64_t> dstShape = {4, 3};
150+ void *rgbDeviceAddr = nullptr;
151+ void *alphaDeviceAddr = nullptr;
152+ void *dstDeviceAddr = nullptr;
153+ aclTensor *rgb = nullptr;
154+ aclTensor *alpha = nullptr;
155+ aclTensor *dst = nullptr;
156+ std::vector<uint8_t> rgbHostData = {10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120};
157+ std::vector<uint8_t> alphaHostData = {255, 255, 255, 255};
158+ std::vector<uint8_t> dstHostData = {1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0};
159+ // 创建rgb aclTensor
160+ ret = CreateAclTensor(rgbHostData, rgbShape, &rgbDeviceAddr, aclDataType::ACL_UINT8, &rgb);
161+ CHECK_RET(ret == ACL_SUCCESS, return ret);
162+ // 创建alpha aclTensor
163+ ret = CreateAclTensor(alphaHostData, alphaShape, &alphaDeviceAddr, aclDataType::ACL_UINT8, &alpha);
164+ CHECK_RET(ret == ACL_SUCCESS, return ret);
165+ // 创建dst aclTensor
166+ ret = CreateAclTensor(dstHostData, dstShape, &dstDeviceAddr, aclDataType::ACL_UINT8, &dst);
167+ CHECK_RET(ret == ACL_SUCCESS, return ret);
168+ 
169+ // 3. 调用CANN算子库API
170+ uint64_t workspaceSize = 0;
171+ aclOpExecutor *executor;
172+ // 调用aclnnMrgba第一段接口
173+ ret = aclnnMrgbaCustomGetWorkspaceSize(rgb, alpha, dst, &workspaceSize, &executor);
174+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMrgbaCustomGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
175+ // 根据第一段接口计算出的workspaceSize申请device内存
176+ void *workspaceAddr = nullptr;
177+ if (workspaceSize > 0) {
178+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
179+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
180+ }
181+ // 调用aclnnMrgba第二段接口
182+ ret = aclnnMrgbaCustom(workspaceAddr, workspaceSize, executor, stream);
183+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMrgbaCustom failed. ERROR: %d\n", ret); return ret);
184+ // 4. (固定写法)同步等待任务执行结束
185+ ret = aclrtSynchronizeStream(stream);
186+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
187+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
188+ auto size = GetShapeSize(dstShape);
189+ std::vector<uint8_t> resultData(size, 0);
190+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), dstDeviceAddr,
191+ size * sizeof(uint8_t),
192+ ACL_MEMCPY_DEVICE_TO_HOST);
193+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
194+ for (int64_t i = 0; i < size; i++) {
195+ LOG_PRINT("result[%ld] is: %u\n", i, resultData[i]);
196+ }
197+ 
198+ // 6. 释放aclTensor
199+ aclDestroyTensor(rgb);
200+ aclDestroyTensor(alpha);
201+ aclDestroyTensor(dst);
202+ 
203+ // 7. 释放device资源，需要根据具体API的接口定义修改
204+ aclrtFree(rgbDeviceAddr);
205+ aclrtFree(alphaDeviceAddr);
206+ aclrtFree(dstDeviceAddr);
207+ if(workspaceSize > 0){
208+ aclrtFree(workspaceAddr);
209+ }
210+ aclrtDestroyStream(stream);
211+ aclrtResetDevice(deviceId);
212+ aclFinalize();
213+ return 0;
214+}
215+ 
216+```
Aobjdetect/mrgba_custom/op_graph/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_op_graph_sources()
Aobjdetect/mrgba_custom/op_graph/mrgba_custom_proto.h+37-0
@@ -0,0 +1,37 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file mgrba_custom_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
16+#define OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+namespace ge
20+{
21+ /**
22+ * @brief Give transparency to the image.
23+ *
24+ * @par Inputs:
25+ * @li rgb: A tensor of the type DT_UINT8.
26+ * @li alpha:A tensor of the type DT_UINT8.
27+ *
28+ * @par Outputs:
29+ * @li dst: A tensor of the type DT_UINT8.
30+ */
31+ REG_OP(MrgbaCustom)
32+ .INPUT(rgb, TensorType({ DT_UINT8 }))
33+ .INPUT(alpha, TensorType({ DT_UINT8 }))
34+ .OUTPUT(dst, TensorType({ DT_UINT8 }))
35+ .OP_END_FACTORY_REG(MrgbaCustom)
36+} // namespace ge
37+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_DETECT_OPS_H_
Aobjdetect/mrgba_custom/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_modules_sources(OPTYPE mrgba_custom ACLNNTYPE aclnn)
Aobjdetect/mrgba_custom/op_host/config/ascend310p/mrgba_custom_binary.json+42-0
@@ -0,0 +1,42 @@
1+{
2+ "op_type": "MrgbaCustom",
3+ "op_list": [
4+ {
5+ "bin_filename": "MrgbaCustom_uint8",
6+ "inputs": [
7+ {
8+ "name": "rgb",
9+ "index": 0,
10+ "dtype": "uint8",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "alpha",
19+ "index": 1,
20+ "dtype": "uint8",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ }
27+ ],
28+ "outputs": [
29+ {
30+ "name": "dst",
31+ "index": 0,
32+ "dtype": "uint8",
33+ "format": "ND",
34+ "paramType": "required",
35+ "shape": [
36+ -2
37+ ]
38+ }
39+ ]
40+ }
41+ ]
42+}
Aobjdetect/mrgba_custom/op_host/config/ascend310p/mrgba_custom_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选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+[MrgbaCustom]
13+default=0
Aobjdetect/mrgba_custom/op_host/mrgba_custom_def.cpp+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/*!
13+ * \file mrgba_custom.cpp
14+ * \brief
15+ */
16+#include "register/op_def_registry.h"
17+ 
18+namespace ops {
19+ class MrgbaCustom : public OpDef {
20+ public:
21+ explicit MrgbaCustom(const char *name) : OpDef(name)
22+ {
23+ this->Input("rgb")
24+ .ParamType(REQUIRED)
25+ .DataType({ge::DT_UINT8})
26+ .Format({ge::FORMAT_ND})
27+ .UnknownShapeFormat({ge::FORMAT_ND});
28+ this->Input("alpha")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_UINT8})
31+ .Format({ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND});
33+ this->Output("dst")
34+ .ParamType(REQUIRED)
35+ .DataType({ge::DT_UINT8})
36+ .Format({ge::FORMAT_ND})
37+ .UnknownShapeFormat({ge::FORMAT_ND});
38+ 
39+ this->AICore().AddConfig("ascend310p");
40+ }
41+ };
42+ 
43+ OP_ADD(MrgbaCustom);
44+}
Aobjdetect/mrgba_custom/op_host/mrgba_custom_infershape.cpp+36-0
@@ -0,0 +1,36 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/*!
13+ * \file mrgba_custom.cc
14+ * \brief
15+ */
16+#include "register/op_impl_registry.h"
17+#include "log/log.h"
18+ 
19+using namespace ge;
20+namespace ops {
21+ static ge::graphStatus InferShape4MrgbaCustom(gert::InferShapeContext* context) {
22+ const gert::Shape *rgb_shape = context->GetInputShape(0);
23+ gert::Shape *dst_shape = context->GetOutputShape(0);
24+ *dst_shape = *rgb_shape;
25+ return GRAPH_SUCCESS;
26+ }
27+ 
28+ static ge::graphStatus InferDataTypeForMrgbaCustom(gert::InferDataTypeContext *context)
29+ {
30+ const ge::DataType dst_dtype = context->GetInputDataType(0);
31+ context->SetOutputDataType(0, dst_dtype);
32+ return GRAPH_SUCCESS;
33+ }
34+ 
35+ IMPL_OP_INFERSHAPE(MrgbaCustom).InferShape(InferShape4MrgbaCustom).InferDataType(InferDataTypeForMrgbaCustom);
36+} // namespace ops
Aobjdetect/mrgba_custom/op_host/mrgba_custom_tiling.cpp+53-0
@@ -0,0 +1,53 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/*!
13+ * \file mrgba_custom_tiling.cc
14+ * \brief
15+ */
16+#include "mrgba_custom_tiling.h"
17+#include "register/op_def_registry.h"
18+#include "log/log.h"
19+ 
20+namespace optiling {
21+ constexpr uint32_t BLOCK_DIM = 8;
22+ 
23+ static ge::graphStatus TilingFuncForMrgbaCustom(gert::TilingContext *context)
24+ {
25+ if (context == nullptr) {
26+ return ge::GRAPH_FAILED;
27+ }
28+ TilingDataMrgba tiling;
29+ auto tensorY = context->GetInputTensor(1);
30+ if (tensorY == nullptr) {
31+ return ge::GRAPH_FAILED;
32+ }
33+ uint32_t totalLength = tensorY->GetShapeSize();
34+ tiling.set_alphaLen(totalLength);
35+ 
36+ context->SetBlockDim(BLOCK_DIM);
37+ if (context->GetRawTilingData() == nullptr) {
38+ return ge::GRAPH_FAILED;
39+ }
40+ tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
41+ context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
42+ return ge::GRAPH_SUCCESS;
43+ }
44+ 
45+ static ge::graphStatus TilingPrepareForMrgbaCustom(gert::TilingParseContext* context) {
46+ OP_LOGD("MrgbaCustom", "TilingPrepareForMrgbaCustom start.");
47+ return ge::GRAPH_SUCCESS;
48+ }
49+ struct MrgbaCustomCompileInfo {};
50+ IMPL_OP_OPTILING(MrgbaCustom)
51+ .Tiling(TilingFuncForMrgbaCustom)
52+ .TilingParse<MrgbaCustomCompileInfo>(TilingPrepareForMrgbaCustom);
53+}
Aobjdetect/mrgba_custom/op_host/mrgba_custom_tiling.h+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file mrgba_custom_tiling.h
13+ * \brief
14+ */
15+#ifndef MRGBA_CUSTOM_TILING_H
16+#define MRGBA_CUSTOM_TILING_H
17+ 
18+#include "register/tilingdata_base.h"
19+ 
20+namespace optiling {
21+ BEGIN_TILING_DATA_DEF(TilingDataMrgba)
22+ TILING_DATA_FIELD_DEF(uint32_t, alphaLen);
23+ END_TILING_DATA_DEF;
24+ 
25+ REGISTER_TILING_DATA_CLASS(MrgbaCustom, TilingDataMrgba)
26+}
27+#endif
Aobjdetect/mrgba_custom/op_kernel/mrgba_custom.cpp+131-0
@@ -0,0 +1,131 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/*!
13+ * \file mrgba_custom.cpp
14+ * \brief
15+ */
16+#include "kernel_operator.h"
17+#include "vector_scheduler.h"
18+ 
19+using namespace AscendC;
20+ 
21+constexpr int32_t BUFFER_NUM = 1;
22+constexpr size_t CHANNEL_RATIO = 3;
23+constexpr float RATIO = 0.003921568627451;// this value is 1/255 for normalize
24+ 
25+class KernelMrgba {
26+public:
27+ __aicore__ inline KernelMrgba()
28+ {
29+ }
30+ 
31+ __aicore__ inline void Init(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR dst, size_t bufferNum, size_t bufferBytes,
32+ size_t gmIdx, size_t gmDataLen)
33+ {
34+ if (bufferBytes <= 0) {
35+ return;
36+ }
37+ pipe.InitBuffer(inQueueRgb, bufferNum, CHANNEL_RATIO * bufferBytes);
38+ pipe.InitBuffer(inQueueAlpha, bufferNum, bufferBytes);
39+ pipe.InitBuffer(outQueueDst, bufferNum, CHANNEL_RATIO * bufferBytes);
40+ 
41+ pipe.InitBuffer(bufAlphaF16C1, bufferBytes * sizeof(half));
42+ pipe.InitBuffer(bufAlphaF16C3, CHANNEL_RATIO * bufferBytes * sizeof(half));
43+ pipe.InitBuffer(bufRgbF16C3, CHANNEL_RATIO * bufferBytes * sizeof(half));
44+ 
45+ rgbGm.SetGlobalBuffer((__gm__ uint8_t *)rgb + CHANNEL_RATIO * gmIdx, CHANNEL_RATIO * gmDataLen);
46+ alphaGm.SetGlobalBuffer((__gm__ uint8_t *)alpha + gmIdx, gmDataLen);
47+ dstGm.SetGlobalBuffer((__gm__ uint8_t *)dst + CHANNEL_RATIO * gmIdx, CHANNEL_RATIO * gmDataLen);
48+ }
49+ 
50+ __aicore__ inline void CalcForAlign32(uint32_t idx, size_t len)
51+ {
52+ if (len <= 0) {
53+ return;
54+ }
55+ uint32_t alphaIdx = idx;
56+ uint32_t rgbIdx = CHANNEL_RATIO * idx;
57+ size_t alphaLen = len;
58+ size_t rgbLen = CHANNEL_RATIO * len;
59+ // copyIn
60+ auto rgbLocal = inQueueRgb.AllocTensor<uint8_t>();
61+ auto alphaLocal = inQueueAlpha.AllocTensor<uint8_t>();
62+ 
63+ DataCopy(alphaLocal, alphaGm[alphaIdx], alphaLen);
64+ DataCopy(rgbLocal, rgbGm[rgbIdx], rgbLen);
65+ inQueueRgb.EnQue(rgbLocal);
66+ inQueueAlpha.EnQue(alphaLocal);
67+ 
68+ // compute
69+ rgbLocal = inQueueRgb.DeQue<uint8_t>();
70+ alphaLocal = inQueueAlpha.DeQue<uint8_t>();
71+ auto dstLocal = outQueueDst.AllocTensor<uint8_t>();
72+ 
73+ auto alphaLocalF16C1 = bufAlphaF16C1.Get<half>();
74+ auto alphaBrbaLocalF16C3 = bufAlphaF16C3.Get<half>();
75+ auto rgbLocalF16C3 = bufRgbF16C3.Get<half>();
76+ 
77+ Cast(alphaLocalF16C1, alphaLocal, RoundMode::CAST_NONE, alphaLen);
78+ 
79+ const uint32_t alphaLocalBrbaShape[2] = {static_cast<uint32_t>(alphaLen), 3};
80+ const uint32_t alphaLocalShape[2] = {static_cast<uint32_t>(alphaLen), 1};
81+ const int32_t broadCastDim = 2;
82+ BroadCast<half, broadCastDim, 1>(alphaBrbaLocalF16C3, alphaLocalF16C1, alphaLocalBrbaShape, alphaLocalShape);
83+ 
84+ half normalizedRatio = RATIO;
85+ Muls(alphaBrbaLocalF16C3, alphaBrbaLocalF16C3, normalizedRatio, rgbLen);
86+ Cast(rgbLocalF16C3, rgbLocal, RoundMode::CAST_NONE, rgbLen);
87+ Mul(rgbLocalF16C3, rgbLocalF16C3, alphaBrbaLocalF16C3, rgbLen);
88+ Cast(dstLocal, rgbLocalF16C3, RoundMode::CAST_FLOOR, rgbLen);
89+ 
90+ outQueueDst.EnQue<uint8_t>(dstLocal);
91+ inQueueRgb.FreeTensor(rgbLocal);
92+ inQueueAlpha.FreeTensor(alphaLocal);
93+ 
94+ // copyOut
95+ dstLocal = outQueueDst.DeQue<uint8_t>();
96+ DataCopy(dstGm[rgbIdx], dstLocal, rgbLen);
97+ outQueueDst.FreeTensor(dstLocal);
98+ }
99+ 
100+protected:
101+ TPipe pipe;
102+ TQue <QuePosition::VECIN, BUFFER_NUM> inQueueRgb;
103+ TQue <QuePosition::VECIN, BUFFER_NUM> inQueueAlpha;
104+ TQue <QuePosition::VECOUT, BUFFER_NUM> outQueueDst;
105+ TBuf <TPosition::VECCALC> bufAlphaF16C1;
106+ TBuf <TPosition::VECCALC> bufAlphaF16C3;
107+ TBuf <TPosition::VECCALC> bufRgbF16C3;
108+ 
109+ GlobalTensor <uint8_t> rgbGm;
110+ GlobalTensor <uint8_t> alphaGm;
111+ GlobalTensor <uint8_t> dstGm;
112+};
113+ 
114+template <typename T>
115+__aicore__ void run_op(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR dst, GM_ADDR tiling, float ubVarNum)
116+{
117+ GET_TILING_DATA(tilingData, tiling);
118+ uint32_t alphaLen = tilingData.alphaLen;
119+ VectorScheduler sch(tilingData.alphaLen, GetBlockNum(), BUFFER_NUM, ubVarNum, sizeof(uint8_t));
120+ KernelMrgba op;
121+ size_t orgVecIdx = GetBlockIdx() * sch.dataLenPerCore;
122+ op.Init(rgb, alpha, dst, sch.bufferNum, sch.dataBytesPerLoop, orgVecIdx, sch.dataLen);
123+ sch.run(&op, sch.dataLen);
124+}
125+ 
126+extern "C" __global__ __aicore__ void mrgba_custom(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR dst, GM_ADDR workspace,
127+ GM_ADDR tiling)
128+{
129+ constexpr float ubVarNum = 100;
130+ run_op<uint8_t>(rgb, alpha, dst, tiling, ubVarNum);
131+}
Aobjdetect/mrgba_custom/op_kernel/vector_scheduler.h+134-0
@@ -0,0 +1,134 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file vector_scheduler.h
13+ * \brief
14+ */
15+#ifndef ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
16+#define ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
17+ 
18+#include "kernel_operator.h"
19+using namespace AscendC;
20+ 
21+constexpr size_t UB_SIZE_BYTE = 248 * 1024;
22+constexpr size_t ALIGN_SIZE_BYTES = 32;
23+constexpr size_t BLOCK_SIZE_BYTES = 32;
24+ 
25+__aicore__ inline size_t UpAlignN(size_t n, size_t N)
26+{
27+ if (N == 0) {
28+ return 0;
29+ }
30+ return (n + N - 1) / N * N;
31+}
32+ 
33+__aicore__ inline size_t DownAlignN(size_t n, size_t N)
34+{
35+ if (N == 0) {
36+ return 0;
37+ }
38+ return n / N * N;
39+}
40+ 
41+__aicore__ inline size_t UpAlign32(size_t n)
42+{
43+ return UpAlignN(n, ALIGN_SIZE_BYTES);
44+}
45+ 
46+__aicore__ inline size_t DownAlign32(size_t n)
47+{
48+ return DownAlignN(n, ALIGN_SIZE_BYTES);
49+}
50+ 
51+class VectorComputer {
52+public:
53+ __aicore__ inline VectorComputer() {};
54+ 
55+ __aicore__ inline void CalcForAlign32(uint32_t idx, size_t len) {};
56+};
57+ 
58+class VectorScheduler {
59+public:
60+ __aicore__ inline VectorScheduler(size_t contentLen, size_t blockDim, size_t bufferNum, float ubVarCount,
61+ size_t sizeofT)
62+ : blockDim(blockDim), bufferNum(bufferNum), ubVarCount(ubVarCount), sizeofT(sizeofT)
63+ {
64+ auto blockIdx = GetBlockIdx();
65+ this->dataLenPer32B = BLOCK_SIZE_BYTES / this->sizeofT;
66+ // L1
67+ this->dataLenPerCore = contentLen / this->blockDim;
68+ if (this->dataLenPerCore < this->dataLenPer32B) {
69+ this->dataLenPerCore = blockIdx == 0 ? contentLen : 0;
70+ this->dataLenTailL1 = 0;
71+ } else {
72+ this->dataLenTailL1 = contentLen % this->blockDim;
73+ }
74+ 
75+ // L2
76+ int maxUbSizePerVar = UB_SIZE_BYTE / ubVarCount / this->bufferNum;
77+ this->dataBytesPerLoop = DownAlign32(maxUbSizePerVar);
78+ this->dataLenPerLoop = this->dataBytesPerLoop / this->sizeofT;
79+ 
80+ this->dataLen = this->dataLenPerCore;
81+ if (blockIdx == this->blockDim - 1) {
82+ this->dataLen += this->dataLenTailL1;
83+ }
84+ this->bufferBytesPerVar = this->dataLen > this->dataLenPerLoop ? this->dataBytesPerLoop : UpAlign32(
85+ this->dataLen * this->sizeofT);
86+ }
87+ 
88+ template<class Computer>
89+ __aicore__ inline void run(Computer *computer, size_t len)
90+ {
91+ if (len <= 0) {
92+ return;
93+ }
94+ size_t loops = len / this->dataLenPerLoop;
95+ size_t tailLen = len % this->dataLenPerLoop;
96+ size_t tailLenA32 = DownAlignN(tailLen, this->dataLenPer32B);
97+ size_t tailLenBackoff = tailLen - tailLenA32;
98+ 
99+ uint32_t idx = 0;
100+ for (size_t i = 0; i < loops; i++) {
101+ computer->CalcForAlign32(idx, this->dataLenPerLoop);
102+ idx = idx + this->dataLenPerLoop;
103+ }
104+ if (tailLenA32) {
105+ idx = loops * this->dataLenPerLoop;
106+ computer->CalcForAlign32(idx, tailLenA32);
107+ }
108+ if (tailLenBackoff > 0) {
109+ idx = len >= this->dataLenPer32B ? len - this->dataLenPer32B : 0;
110+ computer->CalcForAlign32(idx, this->dataLenPer32B);
111+ }
112+ }
113+ 
114+public:
115+ float ubVarCount;
116+ size_t blockDim;
117+ size_t bufferNum;
118+ 
119+ size_t sizeofT;
120+ 
121+ size_t dataLenPer32B;
122+ // L1
123+ size_t dataLen;
124+ size_t dataLenPerCore;
125+ size_t dataLenTailL1;
126+ size_t bufferBytesPerVar;
127+ // L2
128+ size_t dataLenPerLoop;
129+ size_t dataBytesPerLoop;
130+ size_t loopL2;
131+ size_t dataLenTailL2;
132+};
133+ 
134+#endif // ASCENDC_OPERATOR_VECTOR_SCHEDULER_H
Aobjdetect/mrgba_custom/tests/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT OR OP_KERNEL_UT))
13+ list(REMOVE_ITEM CURRENT_DIRS ut)
14+endif()
15+ 
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aobjdetect/mrgba_custom/tests/ut/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+ 
14+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT))
15+ list(REMOVE_ITEM CURRENT_DIRS op_host)
16+endif()
17+ 
18+if(NOT (UT_TEST_ALL OR OP_KERNEL_UT))
19+ list(REMOVE_ITEM CURRENT_DIRS op_kernel)
20+endif()
21+ 
22+foreach(SUB_DIR ${CURRENT_DIRS})
23+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
24+ add_subdirectory(${SUB_DIR})
25+ endif()
26+endforeach()
Aobjdetect/mrgba_custom/tests/ut/op_graph/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/mrgba_custom/tests/ut/op_graph/fusion_pass/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/mrgba_custom/tests/ut/op_host/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+# ----------------------------------------------------------------------------
Aobjdetect/mrgba_custom/tests/ut/op_kernel/CMakeLists.txt+27-0
@@ -0,0 +1,27 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+add_library(mrgba_custom SHARED
12+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel/mrgba_custom.cpp
13+ ./test_mrgba_custom.cpp
14+)
15+ 
16+target_compile_options(mrgba_custom PRIVATE -g -include ${CMAKE_CURRENT_SOURCE_DIR}/test_mrgba_custom.h)
17+ 
18+target_link_libraries(mrgba_custom PRIVATE
19+ $<BUILD_INTERFACE:intf_llt_pub_asan>
20+ -Wl,--whole-archive
21+ -Wl,--no-as-needed
22+ ${PRIVATE_L}
23+ -Wl,--as-needed
24+ -Wl,--no-whole-archive
25+)
26+ 
27+ 
Aobjdetect/mrgba_custom/tests/ut/op_kernel/mrgba_custom_data/gen_data.py+24-0
@@ -0,0 +1,24 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+import sys
13+import numpy as np
14+import torch
15+ 
16+def gen_golden_data_simple():
17+ rgb = np.array([480,640,3]).astype(np.uint8)
18+ alpha = np.array([480,640,1]).astype(np.uint8)
19+ 
20+ rgb.tofile("./rgb.bin")
21+ alpha.tofile("./alpha.bin")
22+ 
23+if __name__ == "__main__":
24+ gen_golden_data_simple()
Aobjdetect/mrgba_custom/tests/ut/op_kernel/test_mrgba_custom.cpp+71-0
@@ -0,0 +1,71 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#include <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+#include "test_mrgba_custom.h"
15+ 
16+#ifdef __CCE_KT_TEST__
17+#include "tikicpulib.h"
18+#include "../data_utils.h"
19+#include "string.h"
20+#include <iostream>
21+#include <string>
22+#endif
23+ 
24+#include <cstdint>
25+ 
26+using namespace std;
27+ 
28+extern "C" __global__ void mrgba_custom(GM_ADDR rgb, GM_ADDR alpha, GM_ADDR dst, GM_ADDR workspace, GM_ADDR tiling);
29+class mrgba_custom_test : public testing::Test {
30+ protected:
31+ static void SetUpTestCase() {
32+ cout << "mrgba_custom_test SetUp\n" << endl;
33+ }
34+ static void TearDownTestCase() {
35+ cout << "mrgba_custom_test TearDown\n" << endl;
36+ }
37+};
38+ 
39+TEST_F(mrgba_custom_test, test_case_uint8) {
40+ // inputs
41+ size_t tiling_data_size = sizeof(MrgbaCustomTilingData);
42+ size_t rgb_size = 480*640*3;
43+ size_t alpha_size = 480*640*1;
44+ uint32_t blockDim = 1;
45+ 
46+ uint8_t *rgb = (uint8_t*)AscendC::GmAlloc(rgb_size);
47+ uint8_t *alpha = (uint8_t*)AscendC::GmAlloc(alpha_size);
48+ uint8_t *dst = (uint8_t*)AscendC::GmAlloc(rgb_size);
49+ 
50+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
51+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
52+ system("cp -r ../../../../../../../ops/objdetect/mrgba_custom/tests/ut/op_kernel/mrgba_custom_data ./");
53+ system("chmod -R 755 ./mrgba_custom_data/");
54+ system("cd ./mrgba_custom_data/ && rm -rf ./*bin");
55+ system("cd ./mrgba_custom_data/ && python3 gen_data.py");
56+ char * path_ = get_current_dir_name();
57+ string path(path_);
58+ ReadFile(path + "/mrgba_custom_data/rgb.bin", rgb_size, rgb, rgb_size);
59+ ReadFile(path + "/mrgba_custom_data/alpha.bin", alpha_size, alpha, alpha_size);
60+ MrgbaCustomTilingData* tilingDatafromBin = reinterpret_cast<MrgbaCustomTilingData*>(tiling);
61+ tilingDatafromBin->alphaLen = 480*640*1;
62+ 
63+ ICPU_RUN_KF(mrgba_custom, blockDim, rgb, alpha, dst, workspace, (uint8_t*)(tilingDatafromBin));
64+ 
65+ AscendC::GmFree(rgb);
66+ AscendC::GmFree(alpha);
67+ AscendC::GmFree(dst);
68+ AscendC::GmFree(workspace);
69+ AscendC::GmFree(tiling);
70+ free(path_);
71+}
Aobjdetect/mrgba_custom/tests/ut/op_kernel/test_mrgba_custom.h+35-0
@@ -0,0 +1,35 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#ifndef _FAST_OP_TEST_MRGBA_CUSTOM_TILING_H_
12+#define _FAST_OP_TEST_MRGBA_CUSTOM_TILING_H_
13+ 
14+#include "kernel_tiling/kernel_tiling.h"
15+ 
16+struct MrgbaCustomTilingData {
17+ uint32_t alphaLen = 0;
18+};
19+ 
20+#pragma pack(1)
21+ 
22+#pragma pack()
23+ 
24+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
25+ __ubuf__ tilingStruct* tilingDataPointer = \
26+ reinterpret_cast<__ubuf__ tilingStruct*>((__ubuf__ uint8_t*)(tilingPointer));
27+ 
28+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
29+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
30+ 
31+#define GET_TILING_DATA(tilingData, tilingPointer) \
32+ MrgbaCustomTilingData tilingData; \
33+ INIT_TILING_DATA(MrgbaCustomTilingData, tilingDataPointer, tilingPointer); \
34+ (tilingData).alphaLen = tilingDataPointer->alphaLen;
35+#endif // _FAST_OP_TEST_MRGBA_CUSTOM_TILING_H_