From 55e6f447cb71afc2bfc8259a3f10f8357dda6168 Mon Sep 17 00:00:00 2001 From: xiao-daipeng Date: Wed, 22 Jun 2022 20:25:47 +0800 Subject: [PATCH 1/7] =?UTF-8?q?=E6=96=B0=E5=A2=9E=E7=AE=97=E5=AD=90uniquec?= =?UTF-8?q?onsecutive=E9=80=82=E9=85=8DV2?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../test_unique_consecutive.py | 65 ++++++++++++++++ third_party/acl/inc/acl/acl_op_compiler.h | 6 ++ third_party/acl/libs/acl_op_compiler.cpp | 16 ++++ .../aten/ops/UniqueConsecutiveKernelNpu.cpp | 74 +++++++++++++++++++ torch_npu/csrc/framework/OpCommand.cpp | 18 ++++- torch_npu/csrc/framework/OpCommand.h | 6 ++ torch_npu/csrc/framework/OpParamMaker.cpp | 70 +++++++++++++----- torch_npu/csrc/framework/OpParamMaker.h | 10 ++- 8 files changed, 240 insertions(+), 25 deletions(-) create mode 100644 test/test_network_ops/test_unique_consecutive.py create mode 100644 torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp diff --git a/test/test_network_ops/test_unique_consecutive.py b/test/test_network_ops/test_unique_consecutive.py new file mode 100644 index 0000000000..e1f4e04077 --- /dev/null +++ b/test/test_network_ops/test_unique_consecutive.py @@ -0,0 +1,65 @@ +# Copyright (c) 2020 Huawei Technologies Co., Ltd +# Copyright (c) 2019, Facebook CORPORATION. +# All rights reserved. +# +# Licensed under the BSD 3-Clause License (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://opensource.org/licenses/BSD-3-Clause +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import torch +import numpy as np +import torch_npu + +from torch_npu.testing.testcase import TestCase, run_tests + +class TestUniqueConsecutive(TestCase): + def test_unique_consecutive(self, device="npu"): + shape_format = [ + [[torch.int32, (2, 3)], 0], + [[torch.long, (2, 3)], 1], + [[torch.float32, (2, 3)], 0], + [[torch.float16, (2, 3)], 1], + [[torch.int32, (2, 3)], None], + [[torch.long, (2, 3)], None], + [[torch.float32, (2, 3)], None], + [[torch.float16, (2, 3)], None] + ] + + for item in shape_format: + cpu_input = torch.rand(item[0][1]).random_(0, 3).to(item[0][0]) + npu_input = cpu_input.npu() + if item[0][0] == torch.float16: + cpu_input = cpu_input.float() + + cpu_output, cpu_idx, cpu_counts = torch.unique_consecutive(cpu_input, return_inverse=True, + return_counts=True, dim=item[1]) + npu_output, npu_idx, npu_counts = torch.unique_consecutive(npu_input, return_inverse=True, + return_counts=True, dim=item[1]) + + if item[0][0] == torch.float16: + cpu_output = cpu_output.half() + self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy()) + self.assertRtolEqual(cpu_idx.numpy(), npu_idx.cpu().numpy()) + self.assertRtolEqual(cpu_counts.numpy(), npu_counts.cpu().numpy()) + + def test_unique_consecutive_case_in_dino(self, device="npu"): + input_list = [ + torch.tensor([224, 224, 96, 96, 96, 96, 96, 96, 96, 96]), + torch.tensor([224, 224]) + ] + for i in input_list: + cpu_output, cpu_counts = torch.unique_consecutive(i, return_counts=True) + npu_output, npu_counts = torch.unique_consecutive(i.npu(), return_counts=True) + self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy()) + self.assertRtolEqual(cpu_counts.numpy(), npu_counts.cpu().numpy()) + +if __name__ == "__main__": + run_tests() \ No newline at end of file diff --git a/third_party/acl/inc/acl/acl_op_compiler.h b/third_party/acl/inc/acl/acl_op_compiler.h index 0807efbd8e..338ec40e4c 100644 --- a/third_party/acl/inc/acl/acl_op_compiler.h +++ b/third_party/acl/inc/acl/acl_op_compiler.h @@ -93,6 +93,12 @@ ACL_FUNC_VISIBILITY aclError aclopCompileAndExecute(const char *opType, const aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag, const char *opPath, aclrtStream stream); +ACL_FUNC_VISIBILITY aclError aclopCompileAndExecuteV2(const char *opType, + int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[], + int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[], + aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag, + const char *opPath, aclrtStream stream); + /** * @ingroup AscendCL * @brief set compile option diff --git a/third_party/acl/libs/acl_op_compiler.cpp b/third_party/acl/libs/acl_op_compiler.cpp index ea5feee968..aaa4886cd7 100644 --- a/third_party/acl/libs/acl_op_compiler.cpp +++ b/third_party/acl/libs/acl_op_compiler.cpp @@ -43,6 +43,22 @@ aclError aclopCompileAndExecute( return 0; } +aclError aclopCompileAndExecuteV2( + const char* opType, + int numInputs, + aclTensorDesc *inputDesc[], + aclDataBuffer *inputs[], + int numOutputs, + aclTensorDesc *outputDesc[], + aclDataBuffer *outputs[], + aclopAttr* attr, + aclopEngineType engineType, + aclopCompileType compileFlag, + const char* opPath, + aclrtStream stream) { + return 0; +} + void GeGeneratorFinalize() { return; } diff --git a/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp b/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp new file mode 100644 index 0000000000..de1eba533a --- /dev/null +++ b/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp @@ -0,0 +1,74 @@ +// Copyright (c) 2020 Huawei Technologies Co., Ltd +// All rights reserved. +// +// Licensed under the BSD 3-Clause License (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://opensource.org/licenses/BSD-3-Clause +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "torch_npu/csrc/framework/utils/OpAdapter.h" +#include "torch_npu/csrc/aten/NPUNativeFunctions.h" + +namespace at_npu { +namespace native { + +std::tuple unique_consecutive_out_npu( + at::Tensor& output, + at::Tensor& inverse_indices, + at::Tensor& counts, + const at::Tensor& self, + const bool return_inverse, + const bool return_counts, + c10::optional dim) { + at::Tensor selfCopy = self; + if (self.scalar_type() == at::ScalarType::Half) { + selfCopy = NPUNativeFunctions::npu_dtype_cast(self, at::ScalarType::Float); + output = NPUNativeFunctions::npu_dtype_cast(output, at::ScalarType::Float); + } + c10::SmallVector output_sync_idx = {0, 2}; + OpCommand cmd; + cmd.Sync(output_sync_idx) + .Name("UniqueConsecutive") + .Input(selfCopy) + .Output(output) + .Output(inverse_indices) + .Output(counts) + .Attr("return_idx", return_inverse) + .Attr("return_counts", return_counts); + if (dim.has_value()) { + cmd.Attr("axis", dim.value()); + } + cmd.Run(); + if (self.scalar_type() == at::ScalarType::Half) { + output = NPUNativeFunctions::npu_dtype_cast(output, at::ScalarType::Half); + } + return std::tie(output, inverse_indices, counts); +} + +std::tuple NPUNativeFunctions::unique_consecutive( + const at::Tensor& self, + const bool return_inverse, + const bool return_counts, + c10::optional dim) { + at::Tensor output = (dim.has_value()) ? + OpPreparation::ApplyTensor(self) : OpPreparation::ApplyTensor(self, {self.numel()}); + at::Tensor inverse_indices = (dim.has_value()) ? + OpPreparation::ApplyTensorWithFormat(self.size(dim.value()), self.options().dtype(at::kLong), ACL_FORMAT_ND) : + OpPreparation::ApplyTensorWithFormat(self.sizes(), self.options().dtype(at::kLong), ACL_FORMAT_ND); + at::Tensor counts = (dim.has_value()) ? + OpPreparation::ApplyTensorWithFormat(self.size(dim.value()), self.options().dtype(at::kLong), ACL_FORMAT_ND) : + OpPreparation::ApplyTensorWithFormat({self.numel()}, self.options().dtype(at::kLong), ACL_FORMAT_ND); + unique_consecutive_out_npu(output, inverse_indices, counts, self, return_inverse, return_counts, dim); + return std::tie(output, inverse_indices, counts); +} + + +} // namespace native +} // namespace at \ No newline at end of file diff --git a/torch_npu/csrc/framework/OpCommand.cpp b/torch_npu/csrc/framework/OpCommand.cpp index 9077b7465e..f818a73c6b 100644 --- a/torch_npu/csrc/framework/OpCommand.cpp +++ b/torch_npu/csrc/framework/OpCommand.cpp @@ -12,6 +12,7 @@ // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. + #include #include "torch_npu/csrc/framework/OpCommand.h" #include "torch_npu/csrc/core/npu/register/OptionsManager.h" @@ -127,26 +128,35 @@ OpCommand& OpCommand::Output( if (!resultTypeDefined && commonType.has_value() && output.scalar_type() != commonType.value()) { output = NPUNativeFunctions::npu_dtype_cast(output, commonType.value()); - } + } ) + outputTensor.emplace_back(output); return AddOutput(output, realType); } void OpCommand::Run() { IF_GRAPH_MODE_THEN_RUN(return;) - if (c10_npu::option::OptionsManager::CheckQueueEnable()) { + if (c10_npu::option::OptionsManager::CheckQueueEnable() && !sync) { ExecuteParas execParams; aclCmd->ExportParams(execParams); c10_npu::queue::QueueParas params(c10_npu::queue::COMPILE_AND_EXECUTE, sizeof(ExecuteParas), &execParams); c10_npu::enCurrentNPUStream(¶ms); aclCmd->releaseSource(false); } else { - aclCmd->Run(); + aclCmd->Run(sync, sync_index, outputTensor); aclCmd->releaseSource(); - } + } aclCmds->Pop(); } +OpCommand& OpCommand::Sync(c10::SmallVector &index) { + sync_index = index; + if (!index.empty()) { + sync = true; + } + return *this; +} + OpCommand& OpCommand::AddTensorInput(at::Tensor &tensor, at::ScalarType forceScaleType, const string &descName, diff --git a/torch_npu/csrc/framework/OpCommand.h b/torch_npu/csrc/framework/OpCommand.h index 2ec73f62cb..a7e6eaa76b 100644 --- a/torch_npu/csrc/framework/OpCommand.h +++ b/torch_npu/csrc/framework/OpCommand.h @@ -117,6 +117,8 @@ public: // Run a single op void Run(); + OpCommand& Sync(c10::SmallVector &sync_index); + private: OpCommand& AddTensorInput(at::Tensor &tensor, at::ScalarType forceScaleType = at::ScalarType::Undefined, @@ -150,6 +152,10 @@ private: c10::optional commonType = c10::nullopt; c10::optional commonShape = c10::nullopt; bool resultTypeDefined = false; + bool sync = false; + c10::SmallVector sync_index; + c10::SmallVector outputTensor; + }; // class OpCommand } // namespace native } // namespace at_npu diff --git a/torch_npu/csrc/framework/OpParamMaker.cpp b/torch_npu/csrc/framework/OpParamMaker.cpp index 5d20035d48..710b6d0508 100644 --- a/torch_npu/csrc/framework/OpParamMaker.cpp +++ b/torch_npu/csrc/framework/OpParamMaker.cpp @@ -107,22 +107,28 @@ namespace at_npu attrValue.data()); } - void OpCommandImpl::Run() - { + void OpCommandImpl::Run( + bool sync, + c10::SmallVector &sync_index, + c10::SmallVector &outputTensor) { NPU_LOGD("Op %s Run.", opName.c_str()); RECORD_FUNCTION(opName, std::vector({})); if (PyGILState_Check()) { // we need to release GIL for NPU to compile op. Py_BEGIN_ALLOW_THREADS - ACL_REQUIRE_OK_OP(InnerRun(opName, execParam), opName.c_str()); + ACL_REQUIRE_OK_OP(InnerRun(opName, execParam, sync, sync_index, outputTensor), opName.c_str()); Py_END_ALLOW_THREADS } else { - ACL_REQUIRE_OK_OP(InnerRun(opName, execParam), opName.c_str()); + ACL_REQUIRE_OK_OP(InnerRun(opName, execParam, sync, sync_index, outputTensor), opName.c_str()); } } - aclError OpCommandImpl::InnerRun(string name, AclExecParam ¶ms) - { + aclError OpCommandImpl::InnerRun( + string name, + AclExecParam ¶ms, + bool sync, + c10::SmallVector &sync_index, + c10::SmallVector &outputTensor) { auto stream = c10_npu::getCurrentNPUStream(); auto inputSize = params.inBuffer.size(); auto outputSize = params.outBuffer.size(); @@ -155,19 +161,45 @@ namespace at_npu TORCH_CHECK(false, "In aoe mode, AclGenGraphAndDumpForOp failed!"); } } - ret = aclopCompileAndExecute( - name.c_str(), - inputSize, - params.inDesc.data(), - params.inBuffer.data(), - outputSize, - params.outDesc.data(), - params.outBuffer.data(), - params.attr, - ACL_ENGINE_SYS, - ACL_COMPILE_SYS, - NULL, - stream); + if (!sync) { + ret = aclopCompileAndExecute( + name.c_str(), + inputSize, + params.inDesc.data(), + params.inBuffer.data(), + outputSize, + params.outDesc.data(), + params.outBuffer.data(), + params.attr, + ACL_ENGINE_SYS, + ACL_COMPILE_SYS, + NULL, + stream); + } else { + int64_t dimSize; + ret = aclopCompileAndExecuteV2( + name.c_str(), + inputSize, + const_cast(params.inDesc.data()), + const_cast(params.inBuffer.data()), + outputSize, + const_cast(params.outDesc.data()), + params.outBuffer.data(), + params.attr, + ACL_ENGINE_SYS, + ACL_COMPILE_SYS, + NULL, + stream); + + for (size_t i = 0; i < sync_index.size(); i++) { + c10::SmallVector real_shape; + for (int64_t j = 0; j < outputTensor[sync_index[i]].dim(); j++) { + C10_NPU_CHECK(aclGetTensorDescDimV2(params.outDesc[sync_index[i]], j, &dimSize)); + real_shape.emplace_back(dimSize); + } + outputTensor[sync_index[i]].resize_(real_shape); + } + } ++index; } while (NpuUtils::IsOomError(ret, index) && (index < NPU_MAX_OP_EXEC_TRY_NUM)); if (reset_flag) diff --git a/torch_npu/csrc/framework/OpParamMaker.h b/torch_npu/csrc/framework/OpParamMaker.h index d1b02a0dea..58b12eba0a 100644 --- a/torch_npu/csrc/framework/OpParamMaker.h +++ b/torch_npu/csrc/framework/OpParamMaker.h @@ -303,7 +303,7 @@ namespace at_npu } } - void Run(); + void Run(bool sync, c10::SmallVector &sync_index, c10::SmallVector &outputTensor); void releaseSource(bool no_blocking = true) { @@ -364,7 +364,13 @@ namespace at_npu } } - aclError InnerRun(string name, AclExecParam ¶ms); + aclError InnerRun( + string name, + AclExecParam ¶ms, + bool sync, + c10::SmallVector &sync_index, + c10::SmallVector &outputTensor + ); private: string opName; -- Gitee From 991272f35575ea127fcd5e5e71d45b59c5f64aba Mon Sep 17 00:00:00 2001 From: pengzirong Date: Mon, 27 Jun 2022 14:50:35 +0800 Subject: [PATCH 2/7] =?UTF-8?q?API=E5=BC=80=E5=8F=91=E6=8C=87=E5=8D=97?= =?UTF-8?q?=E2=80=94=E2=80=94>=E7=AE=97=E5=AD=90=E5=BC=80=E5=8F=91?= =?UTF-8?q?=E6=8C=87=E5=8D=97?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- ...00\345\217\221\346\214\207\345\215\227.md" | 233 +++++++++--------- ...2\351\205\215\346\265\201\347\250\213.png" | Bin .../figures/zh-cn_image_0000001144082048.png" | Bin .../figures/zh-cn_image_0000001144082056.png" | Bin .../figures/zh-cn_image_0000001144082064.png" | Bin .../figures/zh-cn_image_0000001144082072.png" | Bin .../figures/zh-cn_image_0000001144082088.png" | Bin .../figures/zh-cn_image_0000001144241896.png" | Bin .../figures/zh-cn_image_0000001190081791.png" | Bin .../figures/zh-cn_image_0000001190081803.png" | Bin .../figures/zh-cn_image_0000001190201935.png" | Bin .../figures/zh-cn_image_0000001190201951.png" | Bin .../figures/zh-cn_image_0000001190201973.png" | Bin .../public_sys-resources/icon-caution.gif" | Bin .../public_sys-resources/icon-danger.gif" | Bin .../public_sys-resources/icon-note.gif" | Bin .../public_sys-resources/icon-notice.gif" | Bin .../public_sys-resources/icon-tip.gif" | Bin .../public_sys-resources/icon-warning.gif" | Bin 19 files changed, 117 insertions(+), 116 deletions(-) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md" (82%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" (100%) rename "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" => "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" (100%) diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md" similarity index 82% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md" index b1b3aa8e14..aa64a03016 100644 --- "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md" +++ "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md" @@ -1,18 +1,18 @@ -# PyTorch API开发指南 +# PyTorch 算子开发指南 - [简介](#简介md) - [算子开发流程](#算子开发流程md) - [算子开发准备](#算子开发准备md) - [环境准备](#环境准备md) - - [API速查](#API速查md) -- [API适配开发](#API适配开发md) + - [算子速查](#算子速查md) +- [算子适配开发](#算子适配开发md) - [前提条件](#前提条件md) - [获取PyTorch源码](#获取PyTorch源码md) - - [注册API开发](#注册API开发md) + - [注册算子开发](#注册算子开发md) - [概述](#概述md) - - [PyTorch1.8.1 注册API开发](#PyTorch1-8-1-注册API开发md) - - [API适配插件开发](#API适配插件开发md) + - [PyTorch1.8.1 注册算子开发](#PyTorch1-8-1-注册算子开发md) + - [算子适配插件开发](#算子适配插件开发md) - [编译和安装PyTorch框架](#编译和安装PyTorch框架md) -- [API功能验证](#API功能验证md) +- [算子功能验证](#算子功能验证md) - [概述](#概述-0md) - [实现过程](#实现过程md) - [FAQ](#FAQmd) @@ -26,7 +26,7 @@ - [PyTorch编译失败，提示“error: call of overload ....”](#PyTorch编译失败-提示-error-call-of-overloadmd) - [附录](#附录md) - [CMake安装方法](#CMake安装方法md) - - [自定义API导出方法](#自定义API导出方法md) + - [自定义算子导出方法](#自定义算子导出方法md)

简介

### 概述 @@ -35,17 +35,17 @@

算子开发流程

-PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。 +PyTorch算子开发包含TBE算子开发和PyTorch框架下的算子适配。 -1. TBE算子开发：昇腾AI软件栈中不包含相应的算子，需要先完成TBE算子的开发，再进行PyTorch框架下的API适配。 +1. TBE算子开发：昇腾AI软件栈中不包含相应的算子，需要先完成TBE算子的开发，再进行PyTorch框架下的算子适配。 TBE算子开发流程及方法请参见《CANN TBE自定义算子开发指南》。 -2. PyTorch框架下的API适配：昇腾AI软件栈中已实现了相应的TBE算子，可直接进行PyTorch框架适配。 +2. PyTorch框架下的算子适配：昇腾AI软件栈中已实现了相应的TBE算子，可直接进行PyTorch框架适配。 - PyTorch框架下的API适配流程如下所示。 + PyTorch框架下的算子适配流程如下所示。 - **图 1** PyTorch框架下的API适配流程 + **图 1** PyTorch框架下的算子适配流程 ![](figures/PyTorch框架下的算子适配流程.png "PyTorch框架下的算子适配流程") **表 1** 算子开发步骤详解 @@ -68,15 +68,15 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。

准备算子开发及运行验证所依赖的开发环境与运行环境。

-

算子开发准备

+

算子开发准备

2

-

API速查

+

算子速查

-

查看TBE算子支持列表和PyTorchAPI适配列表。

-
  • 当前昇腾AI处理器支持的算子列表及支持的算子的详细规格约束。
  • 当前PyTorch适配的API列表。
+

查看TBE算子支持列表和PyTorch算子适配列表。

+
  • 当前昇腾AI处理器支持的算子列表及支持的算子的详细规格约束。
  • 当前PyTorch适配的算子列表。

3

@@ -85,7 +85,7 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。

获取昇腾社区PyTorch源码。

-

API适配开发

+

算子适配开发

4

@@ -97,9 +97,9 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。

5

-

API适配层开发

+

算子适配层开发

-

API适配层开发，将基于第三方框架的算子属性映射成适配昇腾AI处理器的算子属性。

+

算子适配层开发，将基于第三方框架的算子属性映射成适配昇腾AI处理器的算子属性。

6

@@ -111,11 +111,11 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。

7

-

API功能验证

+

算子功能验证

-

在真实的硬件环境中验证API功能。

+

在真实的硬件环境中验证算子功能。

-

API功能验证

+

算子功能验证

@@ -123,11 +123,12 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。 +

算子开发准备

- **[环境准备](#环境准备md)** -- **[API速查](#API速查md)** +- **[算子速查](#算子速查md)**

环境准备

@@ -153,30 +154,30 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。 -

API速查

+

算子速查

-进行API开发时，您可以查询当前昇腾AI处理器中支持的算子列表和当前PyTorch适配的API列表。根据查询结果进行算子开发或PyTorchAPI适配。 +进行算子开发时，您可以查询当前昇腾AI处理器中支持的算子列表和当前PyTorch适配的算子列表。根据查询结果进行算子开发或PyTorch算子适配。 -- 若昇腾AI处理器不支持算子，则需要进行TBE算子开发和PyTorch框架API适配。 -- 若昇腾AI处理器支持算子但PyTorch框架没有API适配，仅需要进行PyTorch框架下API适配。 -- 若PyTorch框架已经适配API，直接使用API即可，不需要开发和适配。 +- 若昇腾AI处理器不支持算子，则需要进行TBE算子开发和PyTorch框架算子适配。 +- 若昇腾AI处理器支持算子但PyTorch框架没有算子适配，仅需要进行PyTorch框架下算子适配。 +- 若PyTorch框架已经适配算子，直接使用算子即可，不需要开发和适配。 -昇腾AI处理器和PyTorch适配的API查询方式如下。 +昇腾AI处理器和PyTorch适配的算子查询方式如下。 -- 当前昇腾AI处理器中支持的算子以及对应的API约束可以通过以下两种方式查询。 +- 当前昇腾AI处理器中支持的算子以及对应的算子约束可以通过以下两种方式查询。 - 命令行开发方式下，您可以参见《CANN 算子清单 \(Ascend 910\)》进行离线查询。 - MindStudio开发方式下，您可以通过MindStudio进行在线查询，详细查看方法可参见《MindStudio 用户指南》中的“算子&模型速查”章节。 -- 当前PyTorch适配的API列表可以参见[《PyTorch API支持清单》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%20API%E6%94%AF%E6%8C%81%E6%B8%85%E5%8D%95.md)。 +- 当前PyTorch适配的算子列表可以参见[《PyTorch API支持清单》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%20API%E6%94%AF%E6%8C%81%E6%B8%85%E5%8D%95.md)。 -

API适配开发

+

算子适配开发

- **[前提条件](#前提条件md)** - **[获取PyTorch源码](#获取PyTorch源码md)** -- **[注册API开发](#注册API开发md)** +- **[注册算子开发](#注册算子开发md)** -- **[API适配插件开发](#API适配插件开发md)** +- **[算子适配插件开发](#算子适配插件开发md)** - **[编译和安装PyTorch框架](#编译和安装PyTorch框架md)** @@ -188,52 +189,52 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。

获取PyTorch源码

-针对pytorch1.8.1版本，PyTorch源码获取请参见[《PyTorch安装指南》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97.md)中“安装PyTorch框架”章节，完成在"pytorch/pytorch_v1.8.1"目录生成适配昇腾AI处理器的全量代码步骤。将在pytorch/pytorch目录中进行PyTorch API适配开发。 +针对pytorch1.8.1版本，PyTorch源码获取请参见[《PyTorch安装指南》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97.md)中“安装PyTorch框架”章节，完成在"pytorch/pytorch_v1.8.1"目录生成适配昇腾AI处理器的全量代码步骤。将在pytorch/pytorch目录中进行PyTorch 算子适配开发。 -

注册API开发

+

注册算子开发

- **[概述](#概述md)** -- **[PyTorch1.8.1 注册API开发](#PyTorch1-8-1-注册API开发md)** +- **[PyTorch1.8.1 注册算子开发](#PyTorch1-8-1-注册算子开发md)**

概述

-当前制定的NPU适配派发原则是NPUAPI的派发不经过框架公共函数，直接派发成NPU适配的函数，即API执行调用栈中只包含NPU适配的函数调用，不包含框架公共函数。PyTorch框架在编译时，会根据 native\_functions.yaml 的定义，按框架中定义的类型和设备分发原则，生成相应的新API的中间层的调用说明。对于NPU，会生成在 build/aten/src/ATen/NPUType.cpp。 +当前制定的NPU适配派发原则是NPU算子的派发不经过框架公共函数，直接派发成NPU适配的函数，即算子执行调用栈中只包含NPU适配的函数调用，不包含框架公共函数。PyTorch框架在编译时，会根据 native\_functions.yaml 的定义，按框架中定义的类型和设备分发原则，生成相应的新算子的中间层的调用说明。对于NPU，会生成在 build/aten/src/ATen/NPUType.cpp。 -

PyTorch1.8.1 注册API开发

+

PyTorch1.8.1 注册算子开发

-##### 注册API开发方法 +##### 注册算子开发方法 1. 打开native\_functions.yaml文件。 - native\_functions.yaml 文件中，定义了所有API函数原型，包括函数名称和参数等信息，每个API函数支持不同硬件平台的派发信息。该文件所在路径为pytorch/aten/src/ATen/native/native\_functions.yaml。 + native\_functions.yaml 文件中，定义了所有算子函数原型，包括函数名称和参数等信息，每个算子函数支持不同硬件平台的派发信息。该文件所在路径为pytorch/aten/src/ATen/native/native\_functions.yaml。 2. 确定需要派发函数。 - - yaml 中已存在的API + - yaml 中已存在的算子 - 将所有与待适配API相关的函数进行派发。 + 将所有与待适配算子相关的函数进行派发。 - - yaml中未存在的自定义API + - yaml中未存在的自定义算子 - 由于yaml中没有相关API的信息，需要手动添加相关函数，包括函数名，参数信息，返回类型信息。添加规则及方法请参见“pytorch/aten/src/ATen/native/README.md“。 + 由于yaml中没有相关算子的信息，需要手动添加相关函数，包括函数名，参数信息，返回类型信息。添加规则及方法请参见“pytorch/aten/src/ATen/native/README.md“。 ``` - - func：适配API名称(输入参数信息) -> 返回类型 + - func：适配算子名称(输入参数信息) -> 返回类型 ``` ##### 示例 -以torch.add\(\)API为例介绍注册API开发过程。 +以torch.add\(\)算子为例介绍注册算子开发过程。 1. 打开native\_functions.yaml文件。 2. 搜索相关函数。 - 在yaml中搜索add，找到与addAPI相关的函数描述func。由于add是PyTorch内置API，不需要手动添加func。若是自定义API，需要手动添加func。 + 在yaml中搜索add，找到与add算子相关的函数描述func。由于add是PyTorch内置算子，不需要手动添加func。若是自定义算子，需要手动添加func。 -3. 确定API相关函数名称及其类型的func描述。 +3. 确定算子相关函数名称及其类型的func描述。 - add.Tensor 的函数分发描述。 ``` @@ -289,18 +290,18 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。 -

API适配插件开发

+

算子适配插件开发

#### 简介 -用户通过开发API适配插件，实现PyTorch原生API的输入参数、输出参数和属性的格式转换，使转换后的格式与TBE算子的输入参数、输出参数和属性的格式相同。适配昇腾AI处理器的PyTorch源代码中提供了适配关联、类型转换和判别、处理动态shape等相关的方法供用户使用。 +用户通过开发算子适配插件，实现PyTorch原生算子的输入参数、输出参数和属性的格式转换，使转换后的格式与TBE算子的输入参数、输出参数和属性的格式相同。适配昇腾AI处理器的PyTorch源代码中提供了适配关联、类型转换和判别、处理动态shape等相关的方法供用户使用。 #### npu_native_functions.yaml文件介绍 ``` backend: NPU # Backend类型 -cpp_namespace: at_npu::native # 插件中开发API的命名空间 -supported: # 已支持的和PyTorch Native Functions对齐的API +cpp_namespace: at_npu::native # 插件中开发算子的命名空间 +supported: # 已支持的和PyTorch Native Functions对齐的算子 - add.Tensor - add.Scalar - slow_conv3d.out @@ -313,10 +314,10 @@ supported: # 已支持的和PyTorch Native Functions对齐的API - addcdiv_ - addcdiv.out -autograd: # 已支持的和PyTorch Native Functions对齐的继承自Function的具有前反向操作的API +autograd: # 已支持的和PyTorch Native Functions对齐的继承自Function的具有前反向操作的算子 - maxpool2d -custom: # 自定义API，需要提供API格式定义 +custom: # 自定义算子，需要提供算子格式定义 - func: npu_dtype_cast(Tensor self, ScalarType dtype) -> Tensor variants: function, method - func: npu_dtype_cast_(Tensor(a!) self, Tensor src) -> Tensor(a!) @@ -326,21 +327,21 @@ custom: # 自定义API，需要提供API格式定义 - func: npu_get_float_status(Tensor self) -> Tensor variants: function, method -custom_autograd: # 自定义继承自Function的自定义API +custom_autograd: # 自定义继承自Function的自定义算子 - func: npu_convolution(Tensor input, Tensor weight, Tensor? bias, ...) -> Tensor ``` -官方提供的native_functions.yaml文件定义了PyTorch Native Functions的具体API定义和分发细节，在NPU设备上适配官方已定义API，我们不需要重新定义，只需要注册NPU分发即可。由于我们可以根据已支持的API（supported，autograd）对应解析官方yaml文件得到每个函数的具体格式，所以对应的函数声明和注册分发可以自动化完成，API迁移和开发的时候只需要关注对应的实现细节即可。对于自定义API，由于没有具体的API定义，我们需要在npu_native_functions.yaml文件中给出定义，以便对API进行结构化解析从而实现自动化注册和Python接口绑定。 +官方提供的native_functions.yaml文件定义了PyTorch Native Functions的具体算子定义和分发细节，在NPU设备上适配官方已定义算子，我们不需要重新定义，只需要注册NPU分发即可。由于我们可以根据已支持的算子（supported，autograd）对应解析官方yaml文件得到每个函数的具体格式，所以对应的函数声明和注册分发可以自动化完成，算子迁移和开发的时候只需要关注对应的实现细节即可。对于自定义算子，由于没有具体的算子定义，我们需要在npu_native_functions.yaml文件中给出定义，以便对算子进行结构化解析从而实现自动化注册和Python接口绑定。 #### 适配插件实现 -1. 注册API。 +1. 注册算子。 - 根据npu_native_functions.yaml文件介绍，添加相应API信息。 + 根据npu_native_functions.yaml文件介绍，添加相应算子信息。 2. 创建适配插件文件。 - NPU TBE算子适配文件保存在pytorch/torch_npu/csrc/aten/ops目录下，命名风格采用大驼峰，命名格式： + .cpp，如：AddKernelNpu.cpp。 + NPU TBE算子适配文件保存在pytorch/torch_npu/csrc/aten/ops目录下，命名风格采用大驼峰，命名格式：<算子名\> + .cpp，如：AddKernelNpu.cpp。 3. 引入依赖头文件。 @@ -349,25 +350,25 @@ custom_autograd: # 自定义继承自Function的自定义API >![](public_sys-resources/icon-note.gif) **说明：** >工具的功能和使用方法，可查看头文件和源码获得。 -4. 定义API适配主体函数。 +4. 定义算子适配主体函数。 - 根据注册API开发中的分发函数确定自定义API适配主体函数。 + 根据注册算子开发中的分发函数确定自定义算子适配主体函数。 5. 分别实现适配主体函数。 - 实现API适配主题函数，根据TBE算子原型构造得到对应的input、output、attr。 + 实现算子适配主题函数，根据TBE算子原型构造得到对应的input、output、attr。 #### 示例 -以torch.add\(\)API为例介绍API适配开发过程。 +以torch.add\(\)算子为例介绍算子适配开发过程。 -1. 注册API。 +1. 注册算子。 - 将addAPI添加到npu_native_functions.yaml文件的对应位置，用于自动化声明和注册。 + 将add算子添加到npu_native_functions.yaml文件的对应位置，用于自动化声明和注册。 ``` - supported: #已支持的Pytorch Native Functions 对齐的API + supported: #已支持的Pytorch Native Functions 对齐的算子 add.Tensor add_.Tensor add.out @@ -397,7 +398,7 @@ custom_autograd: # 自定义继承自Function的自定义API >" CalcuOpUtil.h "中主要包含类型转换和判别的函数。 >" OpAdapter.h"文件中主要包含适配关联的头文件。 -4. 定义API适配主体函数。 +4. 定义算子适配主体函数。 ``` at::Tensor NPUNativeFunctions::add(const at::Tensor &self, const at::Tensor &other, at::Scalar alpha) @@ -409,7 +410,7 @@ custom_autograd: # 自定义继承自Function的自定义API > ![](public_sys-resources/icon-note.gif) **说明：** > - > NPUNativeFunctions是API定义需要添加的命名空间约束。 + > NPUNativeFunctions是算子定义需要添加的命名空间约束。 5. 分别实现适配主体函数。 @@ -596,7 +597,7 @@ pip3 install --upgrade torch_npu-1.8.1rc1-cp37-cp37m-linux_{arch}.whl 修改代码之后，需要重新执行“编译”和“安装”PyTorch插件过程。 -

API功能验证

+

算子功能验证

- **[概述](#概述-0md)** @@ -607,13 +608,13 @@ pip3 install --upgrade torch_npu-1.8.1rc1-cp37-cp37m-linux_{arch}.whl #### 简介 -完成API适配开发后，可通过运行适配昇腾处理器的PyTorchAPI，验证API运行结果是否正确。 +完成算子适配开发后，可通过运行适配昇腾处理器的PyTorch算子，验证算子运行结果是否正确。 -API功能验证会覆盖API开发的所有交付件，包含实现文件，API原型定义、API信息库以及API适配插件。本节仅对验证的方法做介绍。 +算子功能验证会覆盖算子开发的所有交付件，包含实现文件，算子原型定义、算子信息库以及算子适配插件。本节仅对验证的方法做介绍。 #### 测试用例及测试工具 -进行自定义API功能验证，通过PyTorch前端构造自定义API的函数并运行验证。 +进行自定义算子功能验证，通过PyTorch前端构造自定义算子的函数并运行验证。 在https://gitee.com/ascend/pytorch 中 "pytorch/test/test_network_ops"目录下提供了测试用例及测试工具，供用户参考。 @@ -621,7 +622,7 @@ API功能验证会覆盖API开发的所有交付件，包含实现文件，API #### 简介 -本章通过具体例子，完成PyTorchAPI的功能测试步骤。 +本章通过具体例子，完成PyTorch算子的功能测试步骤。 #### 操作步骤 @@ -634,7 +635,7 @@ API功能验证会覆盖API开发的所有交付件，包含实现文件，API ${HOME}/Ascend/ascend-toolkit/set_env.sh ``` -2. 编写测试脚本。以addAPI为例，在“pytorch/test/test\_network\_ops“路径下编写测试脚本文件： test\_add.py。 +2. 编写测试脚本。以add算子为例，在“pytorch/test/test\_network\_ops“路径下编写测试脚本文件： test\_add.py。 ``` # 引入依赖库 @@ -795,7 +796,7 @@ pip3.7 install torchvision --no-deps #### 现象描述 -完成“自定义TBE算子”开发，和“PyTorch”适配开发，但执行测试用例，发现无法调用到新开发的API。 +完成“自定义TBE算子”开发，和“PyTorch”适配开发，但执行测试用例，发现无法调用到新开发的算子。 #### 可能原因 @@ -805,15 +806,15 @@ pip3.7 install torchvision --no-deps #### 处理方法 -1. 参考“[API功能验证](#API功能验证md)”章节，完成运行环境设置，特别注意： +1. 参考“[算子功能验证](#算子功能验证md)”章节，完成运行环境设置，特别注意： ``` . /home/HwHiAiUser/Ascend/ascend-toolkit/set_env.sh ``` -2. 检查yaml文件中对应API的分发配置，是否正确的完整分发； +2. 检查yaml文件中对应算子的分发配置，是否正确的完整分发； 3. 分析排查代码实现，建议手段： - 1. 修改"pytorch"中的API适配实现，让“test\_add.py”可以调用到“自定义API包”中的TBE算子； + 1. 修改"pytorch"中的算子适配实现，让“test\_add.py”可以调用到“自定义算子包”中的TBE算子； "pytorch/aten/src/ATen/native/npu/AddKernelNpu.cpp" @@ -829,44 +830,44 @@ pip3.7 install torchvision --no-deps 至此步骤，不应该有错误，屏幕应该输出 "add" 中增加的日志打印。若出错，请完成代码清理排查，保证无新开发的代码影响测试。 - 3. 将新开发的“自定义TBE算子”合并到"cann"中，在对应的在API入口增加日志打印，作为运行标识。 + 3. 将新开发的“自定义TBE算子”合并到"cann"中，在对应的在算子入口增加日志打印，作为运行标识。 4. 完成上面的"cann"编译、安装，调用“python3.7.5 test\_add.py”进行测试。 >![](public_sys-resources/icon-note.gif) **说明：** - >根据Ascend的设计逻辑，用户开发安装的“custom”API包优先级高于“built-in”的内置API包，在运行加载时，会优先加载调度“custom”包中的API。过程中，若解析“custom”中的API信息文件失败，则会跳过“custom”API包，不加载调度任何“custom”API包中的任何API。 + >根据Ascend的设计逻辑，用户开发安装的“custom”算子包优先级高于“built-in”的内置算子包，在运行加载时，会优先加载调度“custom”包中的算子。过程中，若解析“custom”中的算子信息文件失败，则会跳过“custom”算子包，不加载调度任何“custom”算子包中的任何算子。 > - >- 若此步骤出错，或屏幕未输出 "add" 中增加的日志打印，则说明新开发的“自定义TBE算子”有错误，影响了“自定义API包”的加载，建议**优先排查新开发的“自定义TBE算子”中的“API信息定义”是否正确**。 - >- 若此步骤正确，至少说明 **新开发的“自定义TBE算子”中的“API信息定义”不影响运行**。 + >- 若此步骤出错，或屏幕未输出 "add" 中增加的日志打印，则说明新开发的“自定义TBE算子”有错误，影响了“自定义算子包”的加载，建议**优先排查新开发的“自定义TBE算子”中的“算子信息定义”是否正确**。 + >- 若此步骤正确，至少说明 **新开发的“自定义TBE算子”中的“算子信息定义”不影响运行**。 5. 调用“python3.7.5 xxx\_testcase.py”进行测试； >![](public_sys-resources/icon-note.gif) **说明：** - >- 若屏幕正常输出新开发的“自定义TBE算子”中增加的日志打印，则至少说明调度到了新开发的API。 + >- 若屏幕正常输出新开发的“自定义TBE算子”中增加的日志打印，则至少说明调度到了新开发的算子。 >- 若屏幕未输出新开发的“自定义TBE算子”中增加的日志打印，则问题可能出现在“PyTorch适配”中，需要排查这一部分的实现代码，较多的可能会出现在“XxxxKernelNpu.cpp”中的输入、输出未能正确适配。

如何确定“TBE算子”是否被“PyTorch适配”正确调用

-不管是“custom”API，还是“built-in”API，在安装后，都是以\*.py源码的方式存放在安装目录下，所以我们可以通过编辑源码，在API入口增加日志的方式，打印输出入参，确定输入的参数是否正确。 +不管是“custom”算子，还是“built-in”算子，在安装后，都是以\*.py源码的方式存放在安装目录下，所以我们可以通过编辑源码，在API入口增加日志的方式，打印输出入参，确定输入的参数是否正确。 >![](public_sys-resources/icon-caution.gif) **注意：** >该操作存在风险，建议在操作前备份计划修改的文件。若未备份，损坏后无法恢复，请及时联系支持人员。 -下面以"built-in"API中的"zn\_2\_nchw"API为例： +下面以"built-in"算子中的"zn\_2\_nchw"算子为例： -1. 打开安装在用户目录下的API包安装目录。 +1. 打开安装在用户目录下的算子包安装目录。 ``` cd ~/.local/Ascend/opp/op_impl/built-in/ai_core/tbe/impl ll ``` - 可以看到相应的API\*.py源码文件均为“只读”，即此时不可编辑。 + 可以看到相应的算子\*.py源码文件均为“只读”，即此时不可编辑。 ![](figures/zh-cn_image_0000001190081791.png) -2. 修改API\*.py源码文件属性，增加“可写”权限。 +2. 修改算子\*.py源码文件属性，增加“可写”权限。 ``` sudo chmod +w zn_2_nchw.py @@ -875,7 +876,7 @@ pip3.7 install torchvision --no-deps ![](figures/zh-cn_image_0000001190081803.png) -3. 打开API\*.py源码文件，增加日志，保存退出。 +3. 打开算子\*.py源码文件，增加日志，保存退出。 ``` vi zn_2_nchw.py @@ -886,8 +887,8 @@ pip3.7 install torchvision --no-deps 上面例子只加了个标识，实际调测时，可以增加打印输入参数信息。 4. 调用脚本执行测试用例，通过打印日志分析“输入参数信息”。 -5. 完成测试分析后，重新打开API\*.py源码文件，删除增加的日志，保存退出。 -6. 修改API\*.py源码文件属性，删除“可写”权限。 +5. 完成测试分析后，重新打开算子\*.py源码文件，删除增加的日志，保存退出。 +6. 修改算子\*.py源码文件属性，删除“可写”权限。 ``` sudo chmod -w zn_2_nchw.py @@ -906,7 +907,7 @@ PyTorch编译失败，提示“ error: ld returned 1 exit status”。 #### 可能原因 -通过日志分析，大概原因为XxxxKernelNpu.cpp中实现的适配函数，与PyTorch框架API要求的分发实现接口参数不匹配。在上面的例子中，是“binary\_cross\_entropy\_npu”，打开对应的XxxxKernelNpu.cpp文件，找到相应的适配函数。 +通过日志分析，大概原因为XxxxKernelNpu.cpp中实现的适配函数，与PyTorch框架算子要求的分发实现接口参数不匹配。在上面的例子中，是“binary\_cross\_entropy\_npu”，打开对应的XxxxKernelNpu.cpp文件，找到相应的适配函数。 ![](figures/zh-cn_image_0000001144241896.png) @@ -944,7 +945,7 @@ PyTorch编译失败，提示“error: call of overload ....”。 - **[CMake安装方法](#CMake安装方法md)** -- **[自定义API导出方法](#自定义API导出方法md)** +- **[自定义算子导出方法](#自定义算子导出方法md)**

CMake安装方法

@@ -990,7 +991,7 @@ PyTorch编译失败，提示“error: call of overload ....”。 #### 简介 -PyTorch训练模型中包含自定义API，将自定义API导出成onnx单API模型，方便转移到其他AI框架中使用。自定义API导出有三种形式，适配NPU的TBE算子导出、C++API导出、纯PythonAPI导出。 +PyTorch训练模型中包含自定义算子，将自定义算子导出成onnx单算子模型，方便转移到其他AI框架中使用。自定义算子导出有三种形式，适配NPU的TBE算子导出、C++算子导出、纯Python算子导出。 #### 前提条件 @@ -1002,10 +1003,10 @@ TBE算子导出有两种方式。 方式一： -1. 定义和注册API +1. 定义和注册算子 ``` - # 定义API + # 定义算子 @parse_args('v', 'v', 'f', 'i', 'i', 'i', 'i') def symbolic_npu_roi_align(g, input, rois, spatial_scale, pooled_height, pooled_width, sample_num, roi_end_mode): args = [input, rois] @@ -1017,7 +1018,7 @@ TBE算子导出有两种方式。 return g.op('torch::npu_roi_align',*args, **kwargs) - # 注册API + # 注册算子 import torch.onnx.symbolic_registry as sym_registry def register_onnx_sym_npu_roi_align(): sym_registry.register_op('npu_roi_align', symbolic_npu_roi_align, '', 11) @@ -1072,7 +1073,7 @@ TBE算子导出有两种方式。 1. 定义方法类 ``` - # 实现API方法类及符号导出实现的方法 + # 实现算子方法类及符号导出实现的方法 class CustomClassOp_Func_npu_roi_align(Function): @staticmethod def forward(ctx, input, rois, spatial_scale, pooled_height, pooled_width , sample_num, roi_end_mode): @@ -1090,10 +1091,10 @@ TBE算子导出有两种方式。 return g.op('torch::npu_roi_align',*args, **kwargs) ``` -2. 自定义API模型 +2. 自定义算子模型 ``` - # 实现API模型 + # 实现算子模型 class NpuOp_npu_roi_align_Module(torch.nn.Module): def __init__(self): super(NpuOp_npu_roi_align_Module, self).__init__() @@ -1151,15 +1152,15 @@ TBE算子导出有两种方式。 >![](public_sys-resources/icon-note.gif) **说明：** >详细实现代码请参见[test\_custom\_ops\_npu\_demo.py](https://gitee.com/ascend/pytorch/blob/master/test/test_npu/test_onnx/torch.onnx/custom_ops_demo/test_custom_ops_npu_demo.py)，如无权限获取代码，请联系华为技术支持申请加入“Ascend”组织。 -#### C++API导出 +#### C++算子导出 -1. 自定义API +1. 自定义算子 ``` import torch import torch_npu import torch.utils.cpp_extension - # 定义C++实现的API + # 定义C++实现的算子 def test_custom_add(): op_source = """ #include @@ -1180,10 +1181,10 @@ TBE算子导出有两种方式。 test_custom_add() ``` -2. 注册自定义API +2. 注册自定义算子 ``` - # 定义API注册方法并注册API + # 定义算子注册方法并注册算子 from torch.onnx import register_custom_op_symbolic def symbolic_custom_add(g, self, other): @@ -1195,16 +1196,16 @@ TBE算子导出有两种方式。 3. 建立模型 ``` - # 建立API模型 + # 建立算子模型 class CustomAddModel(torch.nn.Module): def forward(self, a, b): return torch.ops.custom_namespace.custom_add(a, b) ``` -4. 导出APIonnx模型 +4. 导出算子onnx模型 ``` - # 将API导出onnx模型 + # 将算子导出onnx模型 def do_export(model, inputs, *args, **kwargs): out = torch.onnx._export(model, inputs, "custom_demo.onnx", *args, **kwargs) @@ -1218,9 +1219,9 @@ TBE算子导出有两种方式。 >![](public_sys-resources/icon-note.gif) **说明：** >详细实现代码请参见[test\_custom\_ops\_demo.py](https://gitee.com/ascend/pytorch/blob/master/test/test_npu/test_onnx/torch.onnx/custom_ops_demo/test_custom_ops_demo.py)，如无权限获取代码，请联系华为技术支持申请加入“Ascend”组织。 -#### 纯PythonAPI导出 +#### 纯Python算子导出 -1. 自定义API +1. 自定义算子 ``` import torch @@ -1237,7 +1238,7 @@ TBE算子导出有两种方式。 import math from torch.nn import init - # 定义API类方法 + # 定义算子类方法 class CustomClassOp_Add_F(Function): @staticmethod def forward(ctx, input1,input2): @@ -1257,7 +1258,7 @@ TBE算子导出有两种方式。 2. 建立模型 ``` - # 注册API并建立模型 + # 注册算子并建立模型 class CustomClassOp_Add(torch.nn.Module): def __init__(self): super(CustomClassOp_Add, self).__init__() @@ -1283,7 +1284,7 @@ TBE算子导出有两种方式。 return rtn ``` -3. 导出APIonnx模型 +3. 导出算子onnx模型 ``` ONNX_FILE_NAME = "./custom_python_module_demo.onnx" @@ -1298,7 +1299,7 @@ TBE算子导出有两种方式。 output = model(input) do_export(model, input, opset_version=11) - # 将API导出到onnx模型 + # 将算子导出到onnx模型 test_class_export() ``` diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" similarity index 100% rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" -- Gitee From b84b8d0ec1a360d7100432da968aec646b4c0437 Mon Sep 17 00:00:00 2001 From: wangxiao Date: Mon, 27 Jun 2022 21:43:21 +0800 Subject: [PATCH 3/7] wrap aclopCompileAndExecuteV2 --- torch_npu/csrc/framework/OpParamMaker.cpp | 2 +- .../interface/AclOpCompileInterface.cpp | 21 +++++++++++++++ .../interface/AclOpCompileInterface.h | 27 +++++++++++++++++++ 3 files changed, 49 insertions(+), 1 deletion(-) diff --git a/torch_npu/csrc/framework/OpParamMaker.cpp b/torch_npu/csrc/framework/OpParamMaker.cpp index 710b6d0508..a84e33201f 100644 --- a/torch_npu/csrc/framework/OpParamMaker.cpp +++ b/torch_npu/csrc/framework/OpParamMaker.cpp @@ -177,7 +177,7 @@ namespace at_npu stream); } else { int64_t dimSize; - ret = aclopCompileAndExecuteV2( + ret = AclopCompileAndExecuteV2( name.c_str(), inputSize, const_cast(params.inDesc.data()), diff --git a/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp b/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp index dd464b3a68..c8c67af34a 100644 --- a/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp +++ b/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp @@ -35,6 +35,7 @@ namespace at_npu LOAD_FUNCTION(aclGenGraphAndDumpForOp) LOAD_FUNCTION(aclCreateGraphDumpOpt) LOAD_FUNCTION(aclDestroyGraphDumpOpt) + LOAD_FUNCTION(aclopCompileAndExecuteV2) aclError AclopSetCompileFlag(aclOpCompileFlag flag) { typedef aclError (*aclopSetCompileFlagFunc)(aclOpCompileFlag); @@ -87,5 +88,25 @@ aclError AclDestroyGraphDumpOpt(aclGraphDumpOption* aclGraphDumpOpt) { return func(aclGraphDumpOpt); } +aclError AclopCompileAndExecuteV2(const char *opType, + int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[], + int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[], + aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag, + const char *opPath, aclrtStream stream) { + typedef aclError(*AclopCompileAndExecuteV2Func)(const char *, + int, aclTensorDesc * [], aclDataBuffer * [], + int, aclTensorDesc * [], aclDataBuffer * [], + aclopAttr *, aclopEngineType, aclopCompileType, + const char *, aclrtStream); + static AclopCompileAndExecuteV2Func func = nullptr; + if (func == nullptr) { + func = (AclopCompileAndExecuteV2Func)GET_FUNC(aclopCompileAndExecuteV2); + } + TORCH_CHECK(func, "Failed to find function ", "aclopCompileAndExecuteV2"); + auto ret = func(opType, numInputs, inputDesc, inputs, numOutputs, + outputDesc, outputs, attr, engineType, compileFlag, opPath, stream); + return ret; +} + } // namespace native } // namespace at_npu \ No newline at end of file diff --git a/torch_npu/csrc/framework/interface/AclOpCompileInterface.h b/torch_npu/csrc/framework/interface/AclOpCompileInterface.h index f59f7fe841..85f783dd32 100644 --- a/torch_npu/csrc/framework/interface/AclOpCompileInterface.h +++ b/torch_npu/csrc/framework/interface/AclOpCompileInterface.h @@ -72,6 +72,33 @@ aclGraphDumpOption* AclCreateGraphDumpOpt(); */ aclError AclDestroyGraphDumpOpt(aclGraphDumpOption* aclGraphDumpOpt); +/** + * @ingroup AscendCL + * @brief compile and execute op + * + * @param opType [IN] op type + * @param numInputs [IN] number of inputs + * @param inputDesc [IN] pointer to array of input tensor descriptions + * @param inputs [IN] pointer to array of input buffers + * @param numOutputs [IN] number of outputs + * @param outputDesc [IN] pointer to array of output tensor descriptions + * @param outputs [IN] pointer to array of outputs buffers + * @param attr [IN] pointer to instance of aclopAttr. + * may pass nullptr if the op has no attribute + * @param engineType [IN] engine type + * @param compileFlag [IN] compile flag + * @param opPath [IN] path of op + * @param stream [IN] stream handle + * @retval ACL_ERROR_NONE The function is successfully executed. + * @retval OtherValues Failure + */ +aclError AclopCompileAndExecuteV2(const char *opType, + int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[], + int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[], + aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag, + const char *opPath, aclrtStream stream); + + } // namespace native } // namespace at_npu -- Gitee From 9f3d065e94b9d4b32b04aa1b029a3a6cf8824365 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=9D=8E=E8=B6=85?= Date: Wed, 22 Jun 2022 18:19:23 +0800 Subject: [PATCH 4/7] fix torch.device.type & parse torch.device --- scripts/codegen/gen_python_functions.py | 18 +++++- scripts/codegen/templates/torch_funcs.py | 63 ++++++++++++++++++--- test/test_device.py | 72 ++++++++++++++++++++++++ torch_npu/npu/__init__.py | 3 +- torch_npu/npu/device.py | 3 +- torch_npu/utils/_tensor_str.py | 5 +- torch_npu/utils/device_guard.py | 36 ++++++++++++ torch_npu/utils/module.py | 6 +- torch_npu/utils/tensor_methods.py | 16 +++++- 9 files changed, 204 insertions(+), 18 deletions(-) create mode 100644 test/test_device.py create mode 100644 torch_npu/utils/device_guard.py diff --git a/scripts/codegen/gen_python_functions.py b/scripts/codegen/gen_python_functions.py index 64337fd148..40cd3e8e29 100644 --- a/scripts/codegen/gen_python_functions.py +++ b/scripts/codegen/gen_python_functions.py @@ -218,21 +218,36 @@ def create_python_device_bindings( ) -> None: """Generates Python bindings to ATen functions""" py_device_method_defs: List[str] = [] + device_methods_def_py_dispatch: List[str] = [] grouped = group_filter_overloads(pairs, pred) PY_DEVICE_METHOD_DEF = CodeTemplate("""\ - torch.${name} = torch_npu.${name} + torch.${name} = _${name} +""") + + PY_DEVICE_METHOD_DEF_DISPATCH = CodeTemplate("""\ + +@torch_device_guard +def _${name}(*args, **kwargs): + return torch_npu.${name}(*args, **kwargs) + """) def method_device_def(name): return PY_DEVICE_METHOD_DEF.substitute(name=name) + def method_device_def_dispatch(name): + return PY_DEVICE_METHOD_DEF_DISPATCH.substitute(name=name) + for name in sorted(grouped.keys(), key=lambda x: str(x)): py_device_method_defs.append(method_device_def(name)) + device_methods_def_py_dispatch.append(method_device_def_dispatch(name)) + fm.write_with_template(filename, filename, lambda: { 'device_methods_def_py': py_device_method_defs, + 'device_methods_def_py_dispatch': device_methods_def_py_dispatch }) def query_methods(filepath): @@ -249,7 +264,6 @@ def create_python_device_bindings( raise RuntimeError("In device methods file " + str(fm.install_dir + filename) + " has multi-definition function.") - def load_signatures( native_functions: List[NativeFunction], *, diff --git a/scripts/codegen/templates/torch_funcs.py b/scripts/codegen/templates/torch_funcs.py index 33c9dfe371..d2bb566fed 100644 --- a/scripts/codegen/templates/torch_funcs.py +++ b/scripts/codegen/templates/torch_funcs.py @@ -13,23 +13,70 @@ # See the License for the specific language governing permissions and # limitations under the License. - import torch import torch_npu +from torch_npu.utils.device_guard import torch_device_guard, device + + +@torch_device_guard +def _device(*args, **kwargs): + new_device = torch_npu.new_device(*args, **kwargs) + if new_device.type == torch.npu.native_device: + return device(type=torch_npu.npu.npu_device, index=new_device.index) + return new_device + + +@torch_device_guard +def _tensor(*args, **kwargs): + return torch_npu.tensor(*args, **kwargs) + + +@torch_device_guard +def _full(*args, **kwargs): + return torch_npu.full(*args, **kwargs) + + +@torch_device_guard +def _randint(*args, **kwargs): + return torch_npu.randint(*args, **kwargs) + + +@torch_device_guard +def _range(*args, **kwargs): + return torch_npu.range(*args, **kwargs) + + +@torch_device_guard +def _arange(*args, **kwargs): + return torch_npu.arange(*args, **kwargs) + + +@torch_device_guard +def _empty_with_format(*args, **kwargs): + return torch_npu.empty_with_format(*args, **kwargs) + + +@torch_device_guard +def _npu_dropout_gen_mask(*args, **kwargs): + return torch_npu.npu_dropout_gen_mask(*args, **kwargs) def jit_script(obj, optimize=None, _frames_up=0, _rcb=None): # (Ascend) Disable extension of torch.jit.script return obj + +${device_methods_def_py_dispatch} + def add_torch_funcs(): - torch.device = torch_npu.new_device - torch.tensor = torch_npu.tensor - torch.full = torch_npu.full - torch.randint = torch_npu.randint - torch.range = torch_npu.range - torch.arange = torch_npu.arange - torch.empty_with_format = torch_npu.empty_with_format + torch.device = _device + torch.tensor = _tensor + torch.full = _full + torch.randint = _randint + torch.range = _range + torch.arange = _arange + torch.empty_with_format = _empty_with_format + torch.npu_dropout_gen_mask = _npu_dropout_gen_mask torch.jit.script = jit_script ${device_methods_def_py} \ No newline at end of file diff --git a/test/test_device.py b/test/test_device.py new file mode 100644 index 0000000000..2d4c0f5975 --- /dev/null +++ b/test/test_device.py @@ -0,0 +1,72 @@ +# Copyright (c) 2021, Huawei Technologies.All rights reserved. +# +# Licensed under the BSD 3-Clause License (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://opensource.org/licenses/BSD-3-Clause +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import torch +import torch_npu + +from torch_npu.testing.testcase import TestCase, run_tests + +class TestDevice(TestCase): + def device_monitor(func): + def wrapper(self, *args, **kwargs): + device_id = 0 + torch.npu.set_device(device_id) + npu_device = torch.randn(2).npu(device_id).device + device_types = [ + "npu", + "npu:"+str(device_id), + torch.device("npu:"+str(device_id)), + torch.device("npu:"+str(device_id)).type, + npu_device + ] + for device_type in device_types: + kwargs["device"] = device_type + npu_tensor = func(self, *args, **kwargs) + self.assertEqual(npu_tensor.device.type, "npu") + self.assertEqual(npu_tensor.device.index, device_id) + return wrapper + + + @device_monitor + def test_torch_tensor_to_device(self, device=None): + cpu_tensor = torch.randn(2, 3) + return cpu_tensor.to(device, torch.int64) + + + @device_monitor + def test_torch_tensor_new_empty_with_device_input(self, device=None): + npu_tensor = torch.ones(2, 3).to(device) + return npu_tensor.new_empty((2, 3), dtype=torch.float16, device=device) + + + @device_monitor + def test_torch_func_arange_with_device_input(self, device=None): + return torch.arange(5, dtype=torch.float32, device=device) + + + @device_monitor + def test_torch_func_zeros_with_device_input(self, device=None): + return torch.zeros((2, 3), dtype=torch.int8, device=device) + + + @device_monitor + def test_tensor_method_npu_with_device_input(self, device=None): + if isinstance(device, str): + device = torch.device(device) + cpu_input = torch.randn(2,3) + return cpu_input.npu(device) + + +if __name__ == '__main__': + run_tests() \ No newline at end of file diff --git a/torch_npu/npu/__init__.py b/torch_npu/npu/__init__.py index 957a8d7779..3b6a8dd890 100644 --- a/torch_npu/npu/__init__.py +++ b/torch_npu/npu/__init__.py @@ -15,7 +15,7 @@ # limitations under the License. __all__ = [ - "native_device", "is_initialized", "_lazy_call", "_lazy_init", "init", "set_dump", + "native_device", "npu_device", "is_initialized", "_lazy_call", "_lazy_init", "init", "set_dump", "synchronize", "device_count", "set_device", "current_device", "_get_device_index", "is_available", "device", "device_of", "stream", "current_stream", "default_stream", "init_dump", @@ -37,6 +37,7 @@ __all__ = [ import torch from .device import __device__ as native_device +from .device import __npu_device__ as npu_device from .utils import (is_initialized, _lazy_call, _lazy_init, init, set_dump, synchronize, device_count, set_device, current_device, _get_device_index, is_available, device, device_of, diff --git a/torch_npu/npu/device.py b/torch_npu/npu/device.py index a9607e5f30..fbfa35d822 100644 --- a/torch_npu/npu/device.py +++ b/torch_npu/npu/device.py @@ -13,4 +13,5 @@ # See the License for the specific language governing permissions and # limitations under the License. -__device__ = 'xla' \ No newline at end of file +__device__ = 'xla' +__npu_device__ = 'npu' \ No newline at end of file diff --git a/torch_npu/utils/_tensor_str.py b/torch_npu/utils/_tensor_str.py index f79ee01383..b5abd6629c 100644 --- a/torch_npu/utils/_tensor_str.py +++ b/torch_npu/utils/_tensor_str.py @@ -14,6 +14,7 @@ # limitations under the License. import torch +import torch_npu from torch._tensor_str import _Formatter as SrcFormatter from torch._tensor_str import PRINT_OPTS, _tensor_str_with_formatter, _add_suffixes, get_summarized_data from torch.overrides import has_torch_function_unary, handle_torch_function @@ -60,8 +61,8 @@ def _str_intern(inp): self, tangent = torch.autograd.forward_ad.unpack_dual(inp) if self.device.type != torch._C._get_default_device()\ - or (self.device.type == 'npu' and torch.npu.current_device() != self.device.index): - suffixes.append('device=\'' + str(self.device) + '\'') + or (self.device.type == torch_npu.npu.native_device and torch.npu.current_device() != self.device.index): + suffixes.append('device=\'' + str(torch_npu.npu.npu_device) + ':'+ str(torch.npu.current_device())+'\'') _default_complex_dtype = torch.cdouble if torch.get_default_dtype() == torch.double else torch.cfloat has_default_dtype = self.dtype in (torch.get_default_dtype(), _default_complex_dtype, torch.int64, torch.bool) diff --git a/torch_npu/utils/device_guard.py b/torch_npu/utils/device_guard.py new file mode 100644 index 0000000000..997239c9c6 --- /dev/null +++ b/torch_npu/utils/device_guard.py @@ -0,0 +1,36 @@ +# Copyright (c) 2020 Huawei Technologies Co., Ltd +# All rights reserved. +# +# Licensed under the BSD 3-Clause License (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://opensource.org/licenses/BSD-3-Clause +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import collections +import torch_npu + +device = collections.namedtuple('device', ['type', 'index']) + +def torch_device_guard(func): + # Parse args/kwargs from namedtuple(torch.device) to matched torch.device objects + def wrapper(*args, **kwargs): + if args: + args_list = list(args) + for index, arg in enumerate(args_list): + if isinstance(arg, tuple) and "type='npu'" in str(arg): + args_list[index] = torch_npu.new_device(type=torch_npu.npu.native_device, index=arg.index) + break + args = tuple(args_list) + if kwargs and isinstance(kwargs.get("device"), tuple): + namedtuple_device = kwargs.get("device") + if "type='npu'" in str(namedtuple_device): + kwargs['device'] = torch_npu.new_device(type=torch_npu.npu.native_device, index=namedtuple_device.index) + return func(*args, **kwargs) + return wrapper \ No newline at end of file diff --git a/torch_npu/utils/module.py b/torch_npu/utils/module.py index 063ee5f701..618bb79ad3 100644 --- a/torch_npu/utils/module.py +++ b/torch_npu/utils/module.py @@ -19,6 +19,7 @@ import logging import torch import torch_npu from torch.nn.modules._functions import SyncBatchNorm as sync_batch_norm +from torch_npu.utils.tensor_methods import torch_device_guard def npu(self, device=None): r"""Moves all model parameters and buffers to the npu. @@ -47,7 +48,7 @@ def npu(self, device=None): torch_npu.npu.enable_graph_mode() return self._apply(lambda t: t.npu(device)) - +@torch_device_guard def to(self, *args, **kwargs): if args and isinstance(args[0], str) and 'npu' in args[0]: args = tuple([list(args)[0].replace('npu', torch_npu.npu.native_device)]) @@ -170,7 +171,8 @@ def ddp_forward(self, *inputs, **kwargs): if self.ddp_uneven_inputs_config.ddp_join_enabled: # Notify joined ranks whether they should sync in backwards pass or not. self._check_global_requires_backward_grad_sync(is_joined_rank=False) - if self.device_ids and self.device_type != torch_npu.npu.native_device: + # Note: module.device_type was builded from device.type("npu") inside Class Module + if self.device_ids and self.device_type != torch_npu.npu.npu_device: if len(self.device_ids) == 1: inputs, kwargs = self.to_kwargs(inputs, kwargs, self.device_ids[0]) output = self.module(*inputs[0], **kwargs[0]) diff --git a/torch_npu/utils/tensor_methods.py b/torch_npu/utils/tensor_methods.py index b51611e062..3e647c3ca9 100644 --- a/torch_npu/utils/tensor_methods.py +++ b/torch_npu/utils/tensor_methods.py @@ -17,6 +17,8 @@ import warnings import torch import torch_npu +from torch_npu.utils.device_guard import torch_device_guard, device + warnings.filterwarnings(action="once") warning_str = "The tensor methods of custom operators would cause performance drop." + \ @@ -58,13 +60,12 @@ def npu_confusion_transpose(self, perm, shape, transpose_first): return torch_npu.npu_confusion_transpose(self, perm, shape, transpose_first) +@torch_device_guard def _npu(self, *args, **kwargs): - warnings.warn(warning_str.format("npu")) return torch_npu._C.npu(self, *args, **kwargs) @property def _is_npu(self): - warnings.warn(warning_str.format("is_npu")) return torch_npu._C.is_npu(self) @@ -72,6 +73,7 @@ def _type(self, *args, **kwargs): return torch_npu._C.type(self, *args, **kwargs) +@torch_device_guard def _to(self, *args, **kwargs): return torch_npu._C.to(self, *args, **kwargs) @@ -98,14 +100,23 @@ def _storage(self): return storage_impl(self) +@torch_device_guard def _new_empty(self, *args, **kwargs): return torch_npu._C.new_empty(self, *args, **kwargs) +@torch_device_guard def _new_empty_strided(self, *args, **kwargs): return torch_npu._C.new_empty_strided(self, *args, **kwargs) +@property +def _device(self): + if torch_npu._C.is_npu(self): + return device(type='npu', index=self.get_device()) + return torch.device("cpu") + + def add_tensor_methods(): torch.Tensor.npu_format_cast_ = npu_format_cast_ torch.Tensor.npu_format_cast = npu_format_cast @@ -122,3 +133,4 @@ def add_tensor_methods(): torch.Tensor.storage = _storage torch.Tensor.new_empty = _new_empty torch.Tensor.new_empty_strided = _new_empty_strided + torch.Tensor.device = _device \ No newline at end of file -- Gitee From 54e803545b1254e44065301637d0d48f1879b6e4 Mon Sep 17 00:00:00 2001 From: hxf12345677 Date: Tue, 28 Jun 2022 19:34:34 +0800 Subject: [PATCH 5/7] =?UTF-8?q?=E3=80=90bugfix=E3=80=91=E4=BF=AE=E5=A4=8Da?= =?UTF-8?q?ll=E7=AE=97=E5=AD=90=E8=BE=93=E5=85=A5=E4=B8=BA=E7=A9=BATensor?= =?UTF-8?q?=EF=BC=8C=E8=AE=A1=E7=AE=97=E7=BB=93=E6=9E=9C=E9=94=99=E8=AF=AF?= =?UTF-8?q?=E5=9C=BA=E6=99=AF?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- test/test_network_ops/test_all.py | 13 +++++++++++-- torch_npu/csrc/aten/ops/AllKernelNpu.cpp | 13 ++++++++++++- 2 files changed, 23 insertions(+), 3 deletions(-) diff --git a/test/test_network_ops/test_all.py b/test/test_network_ops/test_all.py index 33bcd3d574..d9de4a6001 100644 --- a/test/test_network_ops/test_all.py +++ b/test/test_network_ops/test_all.py @@ -37,7 +37,7 @@ class TestAll(TestCase): output = output.numpy() return output - def test_all_shape_format(self, device="npu"): + def test_all_shape_format(self): shape_list = [[1024], [32, 1024], [32, 8, 1024], [128, 32, 8, 1024], [2, 0, 2]] for item in shape_list: cpu_input, npu_input = self.create_bool_tensor(item, 0, 1) @@ -71,7 +71,7 @@ class TestAll(TestCase): output1 = output1.to("cpu").numpy() return output0, output1 - def test_alld_shape_format(self, device="npu"): + def test_alld_shape_format(self): shape_list = [[1024], [32, 1024], [32, 8, 1024], [128, 32, 8, 1024]] for item in shape_list: cpu_input, npu_input = self.create_bool_tensor(item, 0, 1) @@ -82,6 +82,15 @@ class TestAll(TestCase): self.assertRtolEqual(cpu_output.astype(np.int32), npu_out0.astype(np.int32)) self.assertRtolEqual(cpu_output.astype(np.int32), npu_out1.astype(np.int32)) + def test_all_tensor_numel_0(self): + ca = torch.rand(1, 2, 0, 3, 4).bool() + na = ca.npu() + cout = ca.all(2) + nout = na.all(2) + cout = cout.numpy() + nout = nout.to("cpu").numpy() + self.assertRtolEqual(cout, nout) + if __name__ == "__main__": run_tests() \ No newline at end of file diff --git a/torch_npu/csrc/aten/ops/AllKernelNpu.cpp b/torch_npu/csrc/aten/ops/AllKernelNpu.cpp index d303b5cddb..d11a357364 100644 --- a/torch_npu/csrc/aten/ops/AllKernelNpu.cpp +++ b/torch_npu/csrc/aten/ops/AllKernelNpu.cpp @@ -62,8 +62,19 @@ at::Tensor& NPUNativeFunctions::all_out( at::Tensor NPUNativeFunctions::all(const at::Tensor& self, int64_t dim, bool keepdim) { TORCH_CHECK(self.scalar_type() == at::ScalarType::Bool || self.scalar_type() == at::ScalarType::Byte, "all only supports torch.uint8 and torch.bool dtypes"); + TORCH_CHECK(dim >= -(self.dim()) && dim < self.dim(), + "The value of dim must be greater than or equal to -self.dim() and less than self.dim()"); if (self.numel() == 0) { - at::Tensor res = OpPreparation::ApplyTensor({}, self.options().dtype(at::kInt), self).fill_(1).to(at::ScalarType::Bool); + c10::SmallVector outputSize; + for(int64_t i = 0; i < self.dim(); i++){ + if(dim != i){ + outputSize.emplace_back(self.size(i)); + } + } + at::Tensor res = OpPreparation::ApplyTensor( + outputSize, + self.options().dtype(at::kInt), + self).fill_(1).to(at::ScalarType::Bool); return res; } -- Gitee From 1022bc44ca2a710ab601a90b1c3f1c825fce55ea Mon Sep 17 00:00:00 2001 From: hxf12345677 Date: Thu, 30 Jun 2022 09:31:24 +0800 Subject: [PATCH 6/7] =?UTF-8?q?bugfix=E3=80=91=20=E4=BF=AE=E5=A4=8Dcdist?= =?UTF-8?q?=E7=AE=97=E5=AD=90=EF=BC=8C=E5=B1=9E=E6=80=A7p=3Dinf,npu?= =?UTF-8?q?=E6=8A=A5=E9=94=99?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- torch_npu/csrc/aten/ops/CdistKernelNpu.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp b/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp index 380c8d63f9..8f03431a8c 100644 --- a/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp +++ b/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp @@ -33,7 +33,7 @@ at::Tensor NPUNativeFunctions::_cdist_forward( // Since double is not supported in NPU, the type of P needs to be converted from double to float. float p_float; if (std::isinf(p)) { - p_float = std::numeric_limits::infinity(); + p_float = -1; } else { TORCH_CHECK(p <= std::numeric_limits::max(), "npu dose not support float64" ); -- Gitee From d68b8c616405b6cb83d5cc752aa4a92454f687ce Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=B0=A4=E5=AE=89=E5=8D=87?= Date: Thu, 30 Jun 2022 14:36:06 +0800 Subject: [PATCH 7/7] Clean code for wrong import order. --- test/test_distributed/run_test.py | 12 ++++++------ torch_npu/csrc/core/npu/NPUGuard.h | 2 +- torch_npu/utils/_tensor_str.py | 5 ++++- 3 files changed, 11 insertions(+), 8 deletions(-) diff --git a/test/test_distributed/run_test.py b/test/test_distributed/run_test.py index 8032a89b6e..2323f4d1f8 100644 --- a/test/test_distributed/run_test.py +++ b/test/test_distributed/run_test.py @@ -1,25 +1,25 @@ #!/usr/bin/env python -from __future__ import print_function import argparse from datetime import datetime -import modulefinder import os import shutil import signal -import subprocess import sys import tempfile import time import unittest import torch -import torch_npu import torch._six -from torch.utils import cpp_extension -from torch.testing._internal.common_utils import TEST_WITH_ROCM, shell import torch.distributed as dist + +from torch.testing._internal.common_utils import TEST_WITH_ROCM, shell + +import torch_npu + + PY2 = sys.version_info <= (3,) PY33 = sys.version_info >= (3, 3) PY36 = sys.version_info >= (3, 6) diff --git a/torch_npu/csrc/core/npu/NPUGuard.h b/torch_npu/csrc/core/npu/NPUGuard.h index 84e81482bf..cbbb774972 100644 --- a/torch_npu/csrc/core/npu/NPUGuard.h +++ b/torch_npu/csrc/core/npu/NPUGuard.h @@ -197,7 +197,7 @@ struct NPUStreamGuard { return guard_.original_device(); } - private: +private: c10::impl::InlineStreamGuard guard_; }; diff --git a/torch_npu/utils/_tensor_str.py b/torch_npu/utils/_tensor_str.py index b5abd6629c..5841116681 100644 --- a/torch_npu/utils/_tensor_str.py +++ b/torch_npu/utils/_tensor_str.py @@ -14,11 +14,14 @@ # limitations under the License. import torch -import torch_npu + from torch._tensor_str import _Formatter as SrcFormatter from torch._tensor_str import PRINT_OPTS, _tensor_str_with_formatter, _add_suffixes, get_summarized_data from torch.overrides import has_torch_function_unary, handle_torch_function +import torch_npu + + class _Formatter(SrcFormatter): def __init__(self, tensor): self.floating_dtype = tensor.dtype.is_floating_point -- Gitee