已合并
chore: 整理 BatchNormGradV3 文档、许可证头与格式 #12098
qianzehong创建于 14 天前
chore: 整理 BatchNormGradV3 文档、许可证头与格式 #12098
已合并
共 23 个文件变更+379-146
| @@ -3486,7 +3486,7 @@ | |||
| 3486 | <td>✓</td> | 3486 | <td>✓</td> |
| 3487 | <td>✓</td> | 3487 | <td>✓</td> |
| 3488 | <td>AI Core</td> | 3488 | <td>AI Core</td> |
| 3489 | - <td>该算子暂无Ascend C代码实现,欢迎开发者补充贡献,贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td> | 3489 | + <td>批归一化反向传播算子,已提供 Ascend C 实现及 ACLNN 接口。</td> |
| 3490 | </tr> | 3490 | </tr> |
| 3491 | <tr> | 3491 | <tr> |
| 3492 | <td>norm</td> | 3492 | <td>norm</td> |
| @@ -1,10 +1,10 @@ | |||
| 1 | # | 1 | # |
| 2 | -# This program is free software, you can redistribute it and/or modify. | ||
| 3 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | -# This file is a part of the CANN Open Software. | 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 8 | # See LICENSE in the root of the software repository for the full text of the License. |
| 9 | #/ | 9 | #/ |
| 10 | 10 | ||
| @@ -17,4 +17,4 @@ foreach(SUB_DIR ${CURRENT_DIRS}) | |||
| 17 | if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | 17 | if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") |
| 18 | add_subdirectory(${SUB_DIR}) | 18 | add_subdirectory(${SUB_DIR}) |
| 19 | endif() | 19 | endif() |
| 20 | -endforeach() | 20 | +endforeach() |
| @@ -1,3 +1,5 @@ | |||
| 1 | # BatchNormGradV3 | 1 | # BatchNormGradV3 |
| 2 | 2 | ||
| 3 | -本目录仅包含BatchNormGradV3算子对应的aclnn接口;如您想要贡献该算子的AscendC实现,请参考[贡献流程](../../CONTRIBUTING.md)。 | 3 | +本目录包含 BatchNormGradV3 算子的 Ascend C Kernel 实现、Host 侧 tiling 与形状和数据类型推导、ACLNN 接口及图模式注册代码,用于批归一化反向传播。 |
| 4 | + | ||
| 5 | +接口说明见 [aclnnBatchNormBackward](docs/aclnnBatchNormBackward.md) 和 [aclnnFastBatchNormBackward](docs/aclnnFastBatchNormBackward.md)。 | ||
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | #include <iostream> | 10 | #include <iostream> |
| 12 | #include <vector> | 11 | #include <vector> |
| @@ -1,11 +1,11 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | 9 | ||
| 10 | message(STATUS "=== Debug: start ops.norm.batch_norm_grad_v3.graph_plugin.CMakeLists.txt ") | 10 | message(STATUS "=== Debug: start ops.norm.batch_norm_grad_v3.graph_plugin.CMakeLists.txt ") |
| 11 | -add_graph_plugin_sources() | 11 | +add_graph_plugin_sources() |
| @@ -1,12 +1,12 @@ | |||
| 1 | # | 1 | # |
| 2 | -# This program is free software, you can redistribute it and/or modify. | ||
| 3 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | -# This file is a part of the CANN Open Software. | 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 8 | # See LICENSE in the root of the software repository for the full text of the License. |
| 9 | #/ | 9 | #/ |
| 10 | message(STATUS "=== Debug: start ops.norm.batch_norm_grad_v3.op_host.CMakeLists.txt ") | 10 | message(STATUS "=== Debug: start ops.norm.batch_norm_grad_v3.op_host.CMakeLists.txt ") |
| 11 | 11 | ||
| 12 | -add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE batch_norm_grad_v3 ACLNNTYPE aclnn_exclude DEPENDENCIES norm_common) | 12 | +add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE batch_norm_grad_v3 ACLNNTYPE aclnn_exclude DEPENDENCIES norm_common) |
| @@ -159,4 +159,4 @@ protected: | |||
| 159 | 159 | ||
| 160 | } // namespace optiling | 160 | } // namespace optiling |
| 161 | 161 | ||
| 162 | -#endif | 162 | +#endif |
| @@ -112,4 +112,4 @@ extern "C" __global__ __aicore__ void batch_norm_grad_v3(GM_ADDR dy, GM_ADDR x, | |||
| 112 | op.Init(gmStruct, &pipe); | 112 | op.Init(gmStruct, &pipe); |
| 113 | op.Process(); | 113 | op.Process(); |
| 114 | } | 114 | } |
| 115 | -} | 115 | +} |
| @@ -1,9 +1,9 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | 9 | ||
| @@ -1,12 +1,12 @@ | |||
| 1 | #!/usr/bin/env python3 | 1 | #!/usr/bin/env python3 |
| 2 | # -*- coding: utf-8 -*- | 2 | # -*- coding: utf-8 -*- |
| 3 | # ---------------------------------------------------------------------------- | 3 | # ---------------------------------------------------------------------------- |
| 4 | -# This program is free software, you can redistribute it and/or modify. | ||
| 5 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 4 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 6 | -# This file is a part of the CANN Open Software. | 5 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 6 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 7 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 9 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 10 | # See LICENSE in the root of the software repository for the full text of the License. |
| 11 | # ---------------------------------------------------------------------------- | 11 | # ---------------------------------------------------------------------------- |
| 12 | 12 | ||
| @@ -16,14 +16,11 @@ from atk.configs.results_config import TaskResult | |||
| 16 | from atk.tasks.api_execute import register | 16 | from atk.tasks.api_execute import register |
| 17 | from atk.tasks.api_execute.base_api import BaseApi | 17 | from atk.tasks.api_execute.base_api import BaseApi |
| 18 | from atk.tasks.api_execute.aclnn_base_api import AclnnBaseApi | 18 | from atk.tasks.api_execute.aclnn_base_api import AclnnBaseApi |
| 19 | -from atk.tasks.dataset.base_dataset import OpsDataset | ||
| 20 | from atk.tasks.backends.lib_interface.acl_wrapper import AclFormat | 19 | from atk.tasks.backends.lib_interface.acl_wrapper import AclFormat |
| 21 | -import numpy as np | ||
| 22 | 20 | ||
| 23 | 21 | ||
| 24 | @register("function_aclnnBatchNormBackward") | 22 | @register("function_aclnnBatchNormBackward") |
| 25 | class aclnnBatchNormBackwardExecutor(BaseApi): | 23 | class aclnnBatchNormBackwardExecutor(BaseApi): |
| 26 | - | ||
| 27 | def __call__(self, input_data: InputDataset, with_output: bool = False): | 24 | def __call__(self, input_data: InputDataset, with_output: bool = False): |
| 28 | gradOut = input_data.kwargs["input0"] | 25 | gradOut = input_data.kwargs["input0"] |
| 29 | input_x = input_data.kwargs["input1"] | 26 | input_x = input_data.kwargs["input1"] |
| @@ -49,7 +46,18 @@ class aclnnBatchNormBackwardExecutor(BaseApi): | |||
| 49 | gradOut_t = gradOut | 46 | gradOut_t = gradOut |
| 50 | input_x_t = input_x | 47 | input_x_t = input_x |
| 51 | 48 | ||
| 52 | - res = torch.ops.aten.native_batch_norm_backward(gradOut_t, input_x_t, weight, runningMean, runningVar, saveMean, saveInvstd, training, eps, outputMask) | 49 | + res = torch.ops.aten.native_batch_norm_backward( |
| 50 | + gradOut_t, | ||
| 51 | + input_x_t, | ||
| 52 | + weight, | ||
| 53 | + runningMean, | ||
| 54 | + runningVar, | ||
| 55 | + saveMean, | ||
| 56 | + saveInvstd, | ||
| 57 | + training, | ||
| 58 | + eps, | ||
| 59 | + outputMask, | ||
| 60 | + ) | ||
| 53 | 61 | ||
| 54 | res_list = [] | 62 | res_list = [] |
| 55 | for i in range(3): | 63 | for i in range(3): |
| @@ -99,4 +107,3 @@ class aclnnBatchNormBackwardAclnnApi(AclnnBaseApi): | |||
| 99 | else: | 107 | else: |
| 100 | output.append(torch.zeros(self.shape_list[i]).to(self.dtype_list[i])) | 108 | output.append(torch.zeros(self.shape_list[i]).to(self.dtype_list[i])) |
| 101 | return output | 109 | return output |
| 102 | - | ||
| @@ -1,9 +1,9 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | 9 | ||
| @@ -14,4 +14,4 @@ foreach(SUB_DIR ${CURRENT_DIRS}) | |||
| 14 | if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | 14 | if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") |
| 15 | add_subdirectory(${SUB_DIR}) | 15 | add_subdirectory(${SUB_DIR}) |
| 16 | endif() | 16 | endif() |
| 17 | -endforeach() | 17 | +endforeach() |
| @@ -1,9 +1,9 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | 9 | ||
| @@ -1,9 +1,9 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | 9 | ||
| @@ -11,4 +11,3 @@ file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOUR | |||
| 11 | if(UT_TEST_ALL OR OP_API_UT) | 11 | if(UT_TEST_ALL OR OP_API_UT) |
| 12 | add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | 12 | add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) |
| 13 | endif() | 13 | endif() |
| 14 | - | ||
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | 10 | ||
| 12 | #include "gtest/gtest.h" | 11 | #include "gtest/gtest.h" |
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | #include <gtest/gtest.h> | 10 | #include <gtest/gtest.h> |
| 12 | #include <iostream> | 11 | #include <iostream> |
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | 10 | ||
| 12 | /*! | 11 | /*! |
| @@ -1,11 +1,11 @@ | |||
| 1 | -# This program is free software, you can redistribute it and/or modify. | ||
| 2 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 1 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 3 | -# This file is a part of the CANN Open Software. | 2 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 3 | +# 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. | 4 | # 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, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 5 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 6 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 7 | # See LICENSE in the root of the software repository for the full text of the License. | 7 | # See LICENSE in the root of the software repository for the full text of the License. |
| 8 | #/ | 8 | #/ |
| 9 | if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE) | 9 | if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE) |
| 10 | AddOpTestCase(batch_norm_grad_v3 "ascend910B1" "" ) | 10 | AddOpTestCase(batch_norm_grad_v3 "ascend910B1" "" ) |
| 11 | -endif() | 11 | +endif() |
| @@ -1,12 +1,12 @@ | |||
| 1 | #!/usr/bin/env python3 | 1 | #!/usr/bin/env python3 |
| 2 | # -*- coding: utf-8 -*- | 2 | # -*- coding: utf-8 -*- |
| 3 | # ---------------------------------------------------------------------------- | 3 | # ---------------------------------------------------------------------------- |
| 4 | -# This program is free software, you can redistribute it and/or modify. | ||
| 5 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 4 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 6 | -# This file is a part of the CANN Open Software. | 5 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 6 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 7 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 9 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 10 | # See LICENSE in the root of the software repository for the full text of the License. |
| 11 | # ---------------------------------------------------------------------------- | 11 | # ---------------------------------------------------------------------------- |
| 12 | import sys | 12 | import sys |
| @@ -14,6 +14,7 @@ import sys | |||
| 14 | import numpy as np | 14 | import numpy as np |
| 15 | from bfloat16 import bfloat16 | 15 | from bfloat16 import bfloat16 |
| 16 | 16 | ||
| 17 | + | ||
| 17 | def compare_data(dx_dtype): | 18 | def compare_data(dx_dtype): |
| 18 | if dx_dtype == "bfloat16": | 19 | if dx_dtype == "bfloat16": |
| 19 | dx = np.fromfile("./cce_cpu_dx.bin", bfloat16).astype(np.float32) | 20 | dx = np.fromfile("./cce_cpu_dx.bin", bfloat16).astype(np.float32) |
| @@ -29,5 +30,6 @@ def compare_data(dx_dtype): | |||
| 29 | if not check_dx: | 30 | if not check_dx: |
| 30 | raise RuntimeError("dx compare failed") | 31 | raise RuntimeError("dx compare failed") |
| 31 | 32 | ||
| 33 | + | ||
| 32 | if __name__ == "__main__": | 34 | if __name__ == "__main__": |
| 33 | - compare_data(*sys.argv[1:]) | 35 | + compare_data(*sys.argv[1:]) |
| @@ -1,12 +1,12 @@ | |||
| 1 | #!/usr/bin/python | 1 | #!/usr/bin/python |
| 2 | # -*- coding: utf-8 -*- | 2 | # -*- coding: utf-8 -*- |
| 3 | # ---------------------------------------------------------------------------- | 3 | # ---------------------------------------------------------------------------- |
| 4 | -# This program is free software, you can redistribute it and/or modify. | ||
| 5 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 4 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 6 | -# This file is a part of the CANN Open Software. | 5 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 6 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 7 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 9 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 10 | # See LICENSE in the root of the software repository for the full text of the License. |
| 11 | # ---------------------------------------------------------------------------- | 11 | # ---------------------------------------------------------------------------- |
| 12 | 12 | ||
| @@ -14,23 +14,48 @@ import sys | |||
| 14 | import numpy as np | 14 | import numpy as np |
| 15 | import torch | 15 | import torch |
| 16 | import tensorflow as tf | 16 | import tensorflow as tf |
| 17 | + | ||
| 17 | bfloat16 = tf.bfloat16.as_numpy_dtype | 18 | bfloat16 = tf.bfloat16.as_numpy_dtype |
| 18 | 19 | ||
| 19 | case_list = { | 20 | case_list = { |
| 20 | - "case1" : {"x_shape":[5, 2, 32, 128], "weight_shape":[2]}, | 21 | + "case1": {"x_shape": [5, 2, 32, 128], "weight_shape": [2]}, |
| 21 | - "case2" : {"x_shape":[170, 2, 2, 125], "weight_shape":[2]}, | 22 | + "case2": {"x_shape": [170, 2, 2, 125], "weight_shape": [2]}, |
| 22 | - "case3" : {"x_shape":[8, 2, 8, 16, 16], "weight_shape":[2]}, | 23 | + "case3": {"x_shape": [8, 2, 8, 16, 16], "weight_shape": [2]}, |
| 23 | - "case4" : {"x_shape":[32, 2, 13, 16], "weight_shape":[8]}, | 24 | + "case4": {"x_shape": [32, 2, 13, 16], "weight_shape": [8]}, |
| 24 | - "case5" : {"x_shape":[32, 2, 16, 16], "weight_shape":[4]}, | 25 | + "case5": {"x_shape": [32, 2, 16, 16], "weight_shape": [4]}, |
| 25 | - "case6" : {"x_shape":[32, 2, 16, 16], "weight_shape":[8]}, | 26 | + "case6": {"x_shape": [32, 2, 16, 16], "weight_shape": [8]}, |
| 26 | - "case7" : {"x_shape":[1, 2, 448, 448], "weight_shape":[2]}, | 27 | + "case7": {"x_shape": [1, 2, 448, 448], "weight_shape": [2]}, |
| 27 | - "case8" : {"x_shape":[10, 2, 192, 108], "weight_shape":[2]}, | 28 | + "case8": {"x_shape": [10, 2, 192, 108], "weight_shape": [2]}, |
| 28 | - "case9" : {"x_shape":[20, 2, 28, 28, 16], "weight_shape":[2]}, | 29 | + "case9": {"x_shape": [20, 2, 28, 28, 16], "weight_shape": [2]}, |
| 29 | - "infer_channel_last_fp32": {"x_shape":[10, 2, 3, 200], "weight_shape":[200], "is_training": False, "epsilon": 0, "format": "NHWC"}, | 30 | + "infer_channel_last_fp32": { |
| 30 | - "infer_fp32": {"x_shape":[1, 2, 1, 1], "weight_shape":[2], "is_training": False, "epsilon": 0, "format": "NCHW"}, | 31 | + "x_shape": [10, 2, 3, 200], |
| 32 | + "weight_shape": [200], | ||
| 33 | + "is_training": False, | ||
| 34 | + "epsilon": 0, | ||
| 35 | + "format": "NHWC", | ||
| 36 | + }, | ||
| 37 | + "infer_fp32": { | ||
| 38 | + "x_shape": [1, 2, 1, 1], | ||
| 39 | + "weight_shape": [2], | ||
| 40 | + "is_training": False, | ||
| 41 | + "epsilon": 0, | ||
| 42 | + "format": "NCHW", | ||
| 43 | + }, | ||
| 31 | } | 44 | } |
| 32 | 45 | ||
| 33 | -def batch_norm_grad_v3(dy, x, weight, running_mean, running_var, save_mean, save_rstd, training=True, epsilon=1e-5, format="NCHW"): | 46 | + |
| 47 | +def batch_norm_grad_v3( | ||
| 48 | + dy, | ||
| 49 | + x, | ||
| 50 | + weight, | ||
| 51 | + running_mean, | ||
| 52 | + running_var, | ||
| 53 | + save_mean, | ||
| 54 | + save_rstd, | ||
| 55 | + training=True, | ||
| 56 | + epsilon=1e-5, | ||
| 57 | + format="NCHW", | ||
| 58 | +): | ||
| 34 | dtype = dy.dtype | 59 | dtype = dy.dtype |
| 35 | if dtype != np.float32: | 60 | if dtype != np.float32: |
| 36 | dy = dy.astype(np.float32) | 61 | dy = dy.astype(np.float32) |
| @@ -47,41 +72,77 @@ def batch_norm_grad_v3(dy, x, weight, running_mean, running_var, save_mean, save | |||
| 47 | tensor_save_rstd = torch.from_numpy(save_rstd).to(torch.float32) | 72 | tensor_save_rstd = torch.from_numpy(save_rstd).to(torch.float32) |
| 48 | 73 | ||
| 49 | if format == "NHWC": | 74 | if format == "NHWC": |
| 50 | - tensor_dy = tensor_dy.permute(0, 3, 1, 2) #NHWC -> NCHW | 75 | + tensor_dy = tensor_dy.permute(0, 3, 1, 2) # NHWC -> NCHW |
| 51 | tensor_x = tensor_x.permute(0, 3, 1, 2) | 76 | tensor_x = tensor_x.permute(0, 3, 1, 2) |
| 52 | elif format == "NDHWC": | 77 | elif format == "NDHWC": |
| 53 | - tensor_dy = tensor_dy.permute(0, 4, 1, 2, 3) #NDHWC -> NCDHW | 78 | + tensor_dy = tensor_dy.permute(0, 4, 1, 2, 3) # NDHWC -> NCDHW |
| 54 | tensor_x = tensor_x.permute(0, 4, 1, 2, 3) | 79 | tensor_x = tensor_x.permute(0, 4, 1, 2, 3) |
| 55 | 80 | ||
| 56 | - res = torch.ops.aten.native_batch_norm_backward(grad_out=tensor_dy, input=tensor_x, | 81 | + res = torch.ops.aten.native_batch_norm_backward( |
| 57 | - weight=tensor_weight, running_mean=tensor_running_mean, running_var=tensor_running_var, | 82 | + grad_out=tensor_dy, |
| 58 | - save_mean=tensor_save_mean, save_invstd=tensor_save_rstd, train=training, eps=epsilon, | 83 | + input=tensor_x, |
| 59 | - output_mask=[True, True, True]) | 84 | + weight=tensor_weight, |
| 85 | + running_mean=tensor_running_mean, | ||
| 86 | + running_var=tensor_running_var, | ||
| 87 | + save_mean=tensor_save_mean, | ||
| 88 | + save_invstd=tensor_save_rstd, | ||
| 89 | + train=training, | ||
| 90 | + eps=epsilon, | ||
| 91 | + output_mask=[True, True, True], | ||
| 92 | + ) | ||
| 60 | 93 | ||
| 61 | dx = res[0].numpy().astype(dtype) | 94 | dx = res[0].numpy().astype(dtype) |
| 62 | dgamma = res[1].numpy().astype(dtype) | 95 | dgamma = res[1].numpy().astype(dtype) |
| 63 | dbeta = res[2].numpy().astype(dtype) | 96 | dbeta = res[2].numpy().astype(dtype) |
| 64 | 97 | ||
| 65 | if format == "NHWC": | 98 | if format == "NHWC": |
| 66 | - dx = np.transpose(dx, (0, 2, 3, 1)) #NCHW -> NHWC | 99 | + dx = np.transpose(dx, (0, 2, 3, 1)) # NCHW -> NHWC |
| 67 | elif format == "NDHWC": | 100 | elif format == "NDHWC": |
| 68 | - dx = np.transpose(dx, (0, 2, 3, 4, 1)) #NCDHW -> NDHWC | 101 | + dx = np.transpose(dx, (0, 2, 3, 4, 1)) # NCDHW -> NDHWC |
| 69 | 102 | ||
| 70 | return dx, dgamma, dbeta | 103 | return dx, dgamma, dbeta |
| 71 | 104 | ||
| 72 | 105 | ||
| 73 | def gen_golden_data_simple(param, dtype): | 106 | def gen_golden_data_simple(param, dtype): |
| 74 | - dy = np.random.uniform(-2, 2, param["x_shape"]).astype(dtype).reshape(param["x_shape"]) | 107 | + dy = ( |
| 75 | - x = np.random.uniform(-2, 2, param["x_shape"]).astype(dtype).reshape(param["x_shape"]) | 108 | + np.random.uniform(-2, 2, param["x_shape"]) |
| 76 | - weight = np.random.uniform(-2, 2, param["weight_shape"]).astype(dtype).reshape(param["weight_shape"]) | 109 | + .astype(dtype) |
| 77 | - running_mean = np.random.uniform(-2, 2, param["weight_shape"]).astype(np.float32).reshape(param["weight_shape"]) | 110 | + .reshape(param["x_shape"]) |
| 78 | - running_var = np.random.uniform(1, 1, param["weight_shape"]).astype(np.float32).reshape(param["weight_shape"]) | 111 | + ) |
| 79 | - save_mean = np.random.uniform(-2, 2, param["weight_shape"]).astype(np.float32).reshape(param["weight_shape"]) | 112 | + x = ( |
| 80 | - save_rstd = np.random.uniform(-2, 2, param["weight_shape"]).astype(np.float32).reshape(param["weight_shape"]) | 113 | + np.random.uniform(-2, 2, param["x_shape"]) |
| 114 | + .astype(dtype) | ||
| 115 | + .reshape(param["x_shape"]) | ||
| 116 | + ) | ||
| 117 | + weight = ( | ||
| 118 | + np.random.uniform(-2, 2, param["weight_shape"]) | ||
| 119 | + .astype(dtype) | ||
| 120 | + .reshape(param["weight_shape"]) | ||
| 121 | + ) | ||
| 122 | + running_mean = ( | ||
| 123 | + np.random.uniform(-2, 2, param["weight_shape"]) | ||
| 124 | + .astype(np.float32) | ||
| 125 | + .reshape(param["weight_shape"]) | ||
| 126 | + ) | ||
| 127 | + running_var = ( | ||
| 128 | + np.random.uniform(1, 1, param["weight_shape"]) | ||
| 129 | + .astype(np.float32) | ||
| 130 | + .reshape(param["weight_shape"]) | ||
| 131 | + ) | ||
| 132 | + save_mean = ( | ||
| 133 | + np.random.uniform(-2, 2, param["weight_shape"]) | ||
| 134 | + .astype(np.float32) | ||
| 135 | + .reshape(param["weight_shape"]) | ||
| 136 | + ) | ||
| 137 | + save_rstd = ( | ||
| 138 | + np.random.uniform(-2, 2, param["weight_shape"]) | ||
| 139 | + .astype(np.float32) | ||
| 140 | + .reshape(param["weight_shape"]) | ||
| 141 | + ) | ||
| 81 | 142 | ||
| 82 | - is_training = param.get('is_training', True) | 143 | + is_training = param.get("is_training", True) |
| 83 | - epsilon = param.get('epsilon', 1e-5) | 144 | + epsilon = param.get("epsilon", 1e-5) |
| 84 | - format = param.get('format', "NCHW") | 145 | + format = param.get("format", "NCHW") |
| 85 | 146 | ||
| 86 | dy.tofile("input_dy.bin") | 147 | dy.tofile("input_dy.bin") |
| 87 | x.tofile("input_x.bin") | 148 | x.tofile("input_x.bin") |
| @@ -91,17 +152,28 @@ def gen_golden_data_simple(param, dtype): | |||
| 91 | save_mean.tofile("input_save_mean.bin") | 152 | save_mean.tofile("input_save_mean.bin") |
| 92 | save_rstd.tofile("input_save_rstd.bin") | 153 | save_rstd.tofile("input_save_rstd.bin") |
| 93 | 154 | ||
| 94 | - dx_golden, dgamma_golden, dbeta_golden = batch_norm_grad_v3(dy, x, weight, running_mean, running_var, save_mean, save_rstd, is_training, epsilon, format) | 155 | + dx_golden, dgamma_golden, dbeta_golden = batch_norm_grad_v3( |
| 156 | + dy, | ||
| 157 | + x, | ||
| 158 | + weight, | ||
| 159 | + running_mean, | ||
| 160 | + running_var, | ||
| 161 | + save_mean, | ||
| 162 | + save_rstd, | ||
| 163 | + is_training, | ||
| 164 | + epsilon, | ||
| 165 | + format, | ||
| 166 | + ) | ||
| 95 | dx_golden.tofile("dx_golden.bin") | 167 | dx_golden.tofile("dx_golden.bin") |
| 96 | dgamma_golden.tofile("dgamma_golden.bin") | 168 | dgamma_golden.tofile("dgamma_golden.bin") |
| 97 | dbeta_golden.tofile("dbeta_golden.bin") | 169 | dbeta_golden.tofile("dbeta_golden.bin") |
| 98 | 170 | ||
| 99 | 171 | ||
| 100 | DTYPE_MAP = { | 172 | DTYPE_MAP = { |
| 101 | - "float16":np.float16, | 173 | + "float16": np.float16, |
| 102 | - "float":np.float32, | 174 | + "float": np.float32, |
| 103 | "float32": np.float32, | 175 | "float32": np.float32, |
| 104 | - "bfloat16": bfloat16 | 176 | + "bfloat16": bfloat16, |
| 105 | } | 177 | } |
| 106 | 178 | ||
| 107 | if __name__ == "__main__": | 179 | if __name__ == "__main__": |
| @@ -112,4 +184,4 @@ if __name__ == "__main__": | |||
| 112 | param = case_list.get(case_name) | 184 | param = case_list.get(case_name) |
| 113 | dtype = DTYPE_MAP.get(dtype) | 185 | dtype = DTYPE_MAP.get(dtype) |
| 114 | print("param:", param) | 186 | print("param:", param) |
| 115 | - gen_golden_data_simple(param, dtype) | 187 | + gen_golden_data_simple(param, dtype) |
| @@ -1,12 +1,12 @@ | |||
| 1 | #!/usr/bin/python | 1 | #!/usr/bin/python |
| 2 | # -*- coding: utf-8 -*- | 2 | # -*- coding: utf-8 -*- |
| 3 | # ---------------------------------------------------------------------------- | 3 | # ---------------------------------------------------------------------------- |
| 4 | -# This program is free software, you can redistribute it and/or modify. | ||
| 5 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. | 4 | # Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 6 | -# This file is a part of the CANN Open Software. | 5 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 6 | +# CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 7 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| 9 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | 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. | 10 | # See LICENSE in the root of the software repository for the full text of the License. |
| 11 | # ---------------------------------------------------------------------------- | 11 | # ---------------------------------------------------------------------------- |
| 12 | 12 | ||
| @@ -14,22 +14,179 @@ import numpy as np | |||
| 14 | import sys | 14 | import sys |
| 15 | 15 | ||
| 16 | TILING_CASE_LIST = { | 16 | TILING_CASE_LIST = { |
| 17 | - "case1": [5, 2, 4096, 8192, 2, 0, 1, 2, 8192, 1, 64, 4096, 0, 32, 8192, 8192, 2, 4096, 4096, 4096, 1, 0, 0, 0, 0], | 17 | + "case1": [ |
| 18 | - "case2": [170, 2, 250, 8000, 2, 0, 1, 2, 8192, 1, 64, 4096, 0, 32, 16000, 16000, 2, 10500, 8000, 8000, 1, 0, 0, 0, 0], | 18 | + 5, |
| 19 | - "case3": [8, 2, 2048, 0, 2, 2, 0, 1, 8192, 1, 64, 4096, 0, 64, 12288, 12288, 1, 16384, 6144, 6144, 2, 4096, 6144, 6144, 2], | 19 | + 2, |
| 20 | + 4096, | ||
| 21 | + 8192, | ||
| 22 | + 2, | ||
| 23 | + 0, | ||
| 24 | + 1, | ||
| 25 | + 2, | ||
| 26 | + 8192, | ||
| 27 | + 1, | ||
| 28 | + 64, | ||
| 29 | + 4096, | ||
| 30 | + 0, | ||
| 31 | + 32, | ||
| 32 | + 8192, | ||
| 33 | + 8192, | ||
| 34 | + 2, | ||
| 35 | + 4096, | ||
| 36 | + 4096, | ||
| 37 | + 4096, | ||
| 38 | + 1, | ||
| 39 | + 0, | ||
| 40 | + 0, | ||
| 41 | + 0, | ||
| 42 | + 0, | ||
| 43 | + ], | ||
| 44 | + "case2": [ | ||
| 45 | + 170, | ||
| 46 | + 2, | ||
| 47 | + 250, | ||
| 48 | + 8000, | ||
| 49 | + 2, | ||
| 50 | + 0, | ||
| 51 | + 1, | ||
| 52 | + 2, | ||
| 53 | + 8192, | ||
| 54 | + 1, | ||
| 55 | + 64, | ||
| 56 | + 4096, | ||
| 57 | + 0, | ||
| 58 | + 32, | ||
| 59 | + 16000, | ||
| 60 | + 16000, | ||
| 61 | + 2, | ||
| 62 | + 10500, | ||
| 63 | + 8000, | ||
| 64 | + 8000, | ||
| 65 | + 1, | ||
| 66 | + 0, | ||
| 67 | + 0, | ||
| 68 | + 0, | ||
| 69 | + 0, | ||
| 70 | + ], | ||
| 71 | + "case3": [ | ||
| 72 | + 8, | ||
| 73 | + 2, | ||
| 74 | + 2048, | ||
| 75 | + 0, | ||
| 76 | + 2, | ||
| 77 | + 2, | ||
| 78 | + 0, | ||
| 79 | + 1, | ||
| 80 | + 8192, | ||
| 81 | + 1, | ||
| 82 | + 64, | ||
| 83 | + 4096, | ||
| 84 | + 0, | ||
| 85 | + 64, | ||
| 86 | + 12288, | ||
| 87 | + 12288, | ||
| 88 | + 1, | ||
| 89 | + 16384, | ||
| 90 | + 6144, | ||
| 91 | + 6144, | ||
| 92 | + 2, | ||
| 93 | + 4096, | ||
| 94 | + 6144, | ||
| 95 | + 6144, | ||
| 96 | + 2, | ||
| 97 | + ], | ||
| 20 | "case4": [32, 2, 208, 6656, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], | 98 | "case4": [32, 2, 208, 6656, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], |
| 21 | "case5": [32, 2, 256, 8192, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], | 99 | "case5": [32, 2, 256, 8192, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], |
| 22 | "case6": [32, 2, 256, 8192, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], | 100 | "case6": [32, 2, 256, 8192, 2, 0, 1, 2, 4096, 0, 64, 2, 1, 1, 1, 2], |
| 23 | - "case7": [1, 2, 200704, 0, 2, 2, 0, 1, 8192, 1, 64, 4096, 0, 64, 14880, 14880, 8, 81664, 7440, 7440, 10, 7264, 7440, 7440, 2], | 101 | + "case7": [ |
| 24 | - "case8": [10, 2, 20736, 0, 2, 2, 0, 1, 8192, 1, 64, 4096, 0, 64, 14880, 14880, 1, 5856, 7440, 7440, 0, 5856, 7440, 7440, 2], | 102 | + 1, |
| 25 | - "case9": [20, 2, 12544, 0, 2, 2, 0, 1, 8192, 1, 64, 4096, 0, 64, 12544, 12544, 1, 0, 6272, 6272, 0, 0, 6272, 6272, 2], | 103 | + 2, |
| 104 | + 200704, | ||
| 105 | + 0, | ||
| 106 | + 2, | ||
| 107 | + 2, | ||
| 108 | + 0, | ||
| 109 | + 1, | ||
| 110 | + 8192, | ||
| 111 | + 1, | ||
| 112 | + 64, | ||
| 113 | + 4096, | ||
| 114 | + 0, | ||
| 115 | + 64, | ||
| 116 | + 14880, | ||
| 117 | + 14880, | ||
| 118 | + 8, | ||
| 119 | + 81664, | ||
| 120 | + 7440, | ||
| 121 | + 7440, | ||
| 122 | + 10, | ||
| 123 | + 7264, | ||
| 124 | + 7440, | ||
| 125 | + 7440, | ||
| 126 | + 2, | ||
| 127 | + ], | ||
| 128 | + "case8": [ | ||
| 129 | + 10, | ||
| 130 | + 2, | ||
| 131 | + 20736, | ||
| 132 | + 0, | ||
| 133 | + 2, | ||
| 134 | + 2, | ||
| 135 | + 0, | ||
| 136 | + 1, | ||
| 137 | + 8192, | ||
| 138 | + 1, | ||
| 139 | + 64, | ||
| 140 | + 4096, | ||
| 141 | + 0, | ||
| 142 | + 64, | ||
| 143 | + 14880, | ||
| 144 | + 14880, | ||
| 145 | + 1, | ||
| 146 | + 5856, | ||
| 147 | + 7440, | ||
| 148 | + 7440, | ||
| 149 | + 0, | ||
| 150 | + 5856, | ||
| 151 | + 7440, | ||
| 152 | + 7440, | ||
| 153 | + 2, | ||
| 154 | + ], | ||
| 155 | + "case9": [ | ||
| 156 | + 20, | ||
| 157 | + 2, | ||
| 158 | + 12544, | ||
| 159 | + 0, | ||
| 160 | + 2, | ||
| 161 | + 2, | ||
| 162 | + 0, | ||
| 163 | + 1, | ||
| 164 | + 8192, | ||
| 165 | + 1, | ||
| 166 | + 64, | ||
| 167 | + 4096, | ||
| 168 | + 0, | ||
| 169 | + 64, | ||
| 170 | + 12544, | ||
| 171 | + 12544, | ||
| 172 | + 1, | ||
| 173 | + 0, | ||
| 174 | + 6272, | ||
| 175 | + 6272, | ||
| 176 | + 0, | ||
| 177 | + 0, | ||
| 178 | + 6272, | ||
| 179 | + 6272, | ||
| 180 | + 2, | ||
| 181 | + ], | ||
| 26 | "infer_channel_last_fp32": [2, 1, 2, 200, 2, 1, 192, 8, 184, 60, 60, 0], | 182 | "infer_channel_last_fp32": [2, 1, 2, 200, 2, 1, 192, 8, 184, 60, 60, 0], |
| 27 | "infer_fp32": [1, 1, 1, 1, 2, 1, 1, 1, 1, 1, 1, 2, 2, 64, 1, 0, 0], | 183 | "infer_fp32": [1, 1, 1, 1, 2, 1, 1, 1, 1, 1, 1, 2, 2, 64, 1, 0, 0], |
| 28 | "infer_full_load_fp32": [1, 2, 1, 8, 1, 1, 956, 2, 0], | 184 | "infer_full_load_fp32": [1, 2, 1, 8, 1, 1, 956, 2, 0], |
| 29 | "infer_split_load_fp32": [10000, 2, 100, 1000000, 1, 1, 6128, 1, 2, 3872, 2, 0], | 185 | "infer_split_load_fp32": [10000, 2, 100, 1000000, 1, 1, 6128, 1, 2, 3872, 2, 0], |
| 30 | "train_full_load_fp32": [1, 2, 1, 8, 1, 1, 956, 2, 0], | 186 | "train_full_load_fp32": [1, 2, 1, 8, 1, 1, 956, 2, 0], |
| 31 | - "train_split_load_fp32": [10000, 2, 100, 1000000, 1, 1, 6128, 1, 2, 3872, 2, 0] | 187 | + "train_split_load_fp32": [10000, 2, 100, 1000000, 1, 1, 6128, 1, 2, 3872, 2, 0], |
| 32 | - } | 188 | +} |
| 189 | + | ||
| 33 | 190 | ||
| 34 | def main(case_name): | 191 | def main(case_name): |
| 35 | data = TILING_CASE_LIST.get(case_name) | 192 | data = TILING_CASE_LIST.get(case_name) |
| @@ -39,6 +196,6 @@ def main(case_name): | |||
| 39 | tiling_data.tofile(tiling_file) | 196 | tiling_data.tofile(tiling_file) |
| 40 | 197 | ||
| 41 | 198 | ||
| 42 | -if __name__ == '__main__': | 199 | +if __name__ == "__main__": |
| 43 | case_name = sys.argv[1] | 200 | case_name = sys.argv[1] |
| 44 | main(case_name) | 201 | main(case_name) |
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | 10 | ||
| 12 | /*! | 11 | /*! |
| @@ -259,4 +258,4 @@ inline void InitTilingData(uint8_t* tiling, T* const_data) | |||
| 259 | tiling_struct tiling_data; \ | 258 | tiling_struct tiling_data; \ |
| 260 | InitTilingData<tiling_struct>(tiling_arg, &tiling_data) | 259 | InitTilingData<tiling_struct>(tiling_arg, &tiling_data) |
| 261 | 260 | ||
| 262 | -#endif | 261 | +#endif |
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | 10 | ||
| 12 | /*! | 11 | /*! |