mouliangyu · mouliangyu · May 28, 2026 · May 28, 2026
diff --git a/docs/designs/module-backend-attr-and-mixed-fatobj.md b/docs/designs/module-backend-attr-and-mixed-fatobj.md
diff --git a/include/PTO/IR/PTO.h b/include/PTO/IR/PTO.h
@@ -185,6 +185,9 @@ inline constexpr llvm::StringLiteral kPTOSimtEntryAttrName = "pto.simt_entry";
 /// Return true if the function carries an explicit entry marker.
 bool hasExplicitPTOEntryAttr(func::FuncOp func);
 
+/// Return true if the function is marked as a PTO device kernel.
+bool isPTOKernelFunction(func::FuncOp func);
+
 /// Return true if the function should be emitted as an AICORE entry.
 bool isPTOEntryFunction(func::FuncOp func);
 

diff --git a/lib/PTO/IR/PTO.cpp b/lib/PTO/IR/PTO.cpp
@@ -2111,6 +2111,11 @@ bool mlir::pto::hasExplicitPTOEntryAttr(func::FuncOp func) {
                   func->hasAttrOfType<UnitAttr>(kLegacyHACCEntryAttrName));
 }
 
+bool mlir::pto::isPTOKernelFunction(func::FuncOp func) {
+  return func && (func->hasAttrOfType<UnitAttr>("pto.kernel") ||
+                  func->hasAttrOfType<UnitAttr>("pto.aicore"));
+}
+
 static constexpr StringLiteral kEffectivePTOEntryAttrName =
     "pto.internal.entry";
 

diff --git a/lib/PTO/Transforms/PTOToEmitC.cpp b/lib/PTO/Transforms/PTOToEmitC.cpp
@@ -2740,14 +2740,18 @@ struct FuncToEmitC : public OpConversionPattern<func::FuncOp> {
     }
 
     if (op.isDeclaration()) {
-      emitcFunc.setSpecifiersAttr(rewriter.getStrArrayAttr({"extern"}));
+      if (op->hasAttr("pto.external_abi"))
+        emitcFunc.setSpecifiersAttr(
+            rewriter.getStrArrayAttr({"extern \"C\"", "AICORE"}));
+      else
+        emitcFunc.setSpecifiersAttr(rewriter.getStrArrayAttr({"extern"}));
       rewriter.eraseOp(op);
       return success();
     }
 
     if (pto::isPTOEntryFunction(op)) {
       emitcFunc.setSpecifiersAttr(
-          rewriter.getStrArrayAttr({"__global__ AICORE"}));
+          rewriter.getStrArrayAttr({"extern \"C\"", "__global__ AICORE"}));
     } else if (op.isPrivate()) {
       emitcFunc.setSpecifiersAttr(
           rewriter.getStrArrayAttr({"static", "AICORE"}));

diff --git a/lib/PTO/Transforms/VPTOLLVMEmitter.cpp b/lib/PTO/Transforms/VPTOLLVMEmitter.cpp
@@ -55,16 +55,9 @@ void attachHIVMKernelAnnotations(llvm::Module &llvmModule);
 
 namespace {
 
-constexpr llvm::StringLiteral kPTOKernelAttrName = "pto.kernel";
-constexpr llvm::StringLiteral kLegacyPTOAICoreAttrName = "pto.aicore";
 constexpr llvm::StringLiteral kVectorSuffix = "_mix_aiv";
 constexpr llvm::StringLiteral kCubeSuffix = "_mix_aic";
 
-static bool hasVPTOKernelAttr(Operation *op) {
-  return op->hasAttr(kPTOKernelAttrName) ||
-         op->hasAttr(kLegacyPTOAICoreAttrName);
-}
-
 static std::string getElementTypeFragment(Type type);
 static Type getElementTypeFromVectorLike(Type type);
 static std::optional<int64_t> getElementCountFromVectorLike(Type type);
@@ -7784,7 +7777,7 @@ static LogicalResult renameKernelFunctionsForKernelKind(ModuleOp module,
   }
 
   for (func::FuncOp funcOp : module.getOps<func::FuncOp>()) {
-    if (!hasVPTOKernelAttr(funcOp))
+    if (!pto::isPTOKernelFunction(funcOp))
       continue;
     if (funcOp.getSymName().ends_with(suffix))
       continue;

diff --git a/lib/PTO/Transforms/VPTOSplitCVModule.cpp b/lib/PTO/Transforms/VPTOSplitCVModule.cpp
@@ -26,10 +26,6 @@ using namespace mlir::pto;
 
 namespace {
 
-static bool hasVPTOKernelAttr(Operation *op) {
-  return op->hasAttr("pto.kernel") || op->hasAttr("pto.aicore");
-}
-
 static bool hasKernelKind(ModuleOp module) {
   return module->hasAttr(FunctionKernelKindAttr::name);
 }
@@ -42,7 +38,7 @@ static bool hasKernelKindChildModule(ModuleOp module) {
 static bool hasCVSections(ModuleOp module) {
   bool found = false;
   module.walk([&](func::FuncOp funcOp) {
-    if (found || !hasVPTOKernelAttr(funcOp))
+    if (found || !pto::isPTOKernelFunction(funcOp))
       return WalkResult::advance();
     WalkResult result = funcOp.walk([&](Operation *op) {
       if (isa<SectionCubeOp, SectionVectorOp>(op)) {
@@ -60,7 +56,7 @@ static bool hasCVSections(ModuleOp module) {
 static bool hasSectionKind(ModuleOp module, FunctionKernelKind kind) {
   bool found = false;
   module.walk([&](func::FuncOp funcOp) {
-    if (found || !hasVPTOKernelAttr(funcOp))
+    if (found || !pto::isPTOKernelFunction(funcOp))
       return WalkResult::advance();
     WalkResult result = funcOp.walk([&](Operation *op) {
       bool matches = kind == FunctionKernelKind::Cube
@@ -125,7 +121,7 @@ static LogicalResult verifyNoNestedSections(ModuleOp module) {
 static LogicalResult verifyKernelFunctionsUseSections(ModuleOp module) {
   LogicalResult status = success();
   module.walk([&](func::FuncOp funcOp) {
-    if (failed(status) || !hasVPTOKernelAttr(funcOp))
+    if (failed(status) || !pto::isPTOKernelFunction(funcOp))
       return WalkResult::advance();
     if (!hasAnySection(funcOp)) {
       status = funcOp.emitOpError(
@@ -141,7 +137,7 @@ static LogicalResult verifyKernelFunctionsUseSections(ModuleOp module) {
 static LogicalResult verifyUniqueSectionKindsPerFunction(ModuleOp module) {
   LogicalResult status = success();
   module.walk([&](func::FuncOp funcOp) {
-    if (failed(status) || !hasVPTOKernelAttr(funcOp))
+    if (failed(status) || !pto::isPTOKernelFunction(funcOp))
       return WalkResult::advance();
     unsigned cubeCount = 0;
     unsigned vectorCount = 0;
@@ -168,7 +164,7 @@ static void eraseKernelFunctionsWithoutSectionKind(ModuleOp module,
                                                    FunctionKernelKind kind) {
   SmallVector<func::FuncOp> eraseFuncs;
   module.walk([&](func::FuncOp funcOp) {
-    if (hasVPTOKernelAttr(funcOp) && !hasSectionKind(funcOp, kind))
+    if (pto::isPTOKernelFunction(funcOp) && !hasSectionKind(funcOp, kind))
       eraseFuncs.push_back(funcOp);
   });
 

diff --git a/test/lit/vpto/backend_attr_invalid.pto b/test/lit/vpto/backend_attr_invalid.pto
@@ -0,0 +1,21 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: not ptoas --pto-arch=a5 --emit-vpto %s -o - 2>&1 | FileCheck %s
+
+module attributes {
+  pto.backend = "other",
+  pto.kernel_kind = #pto.kernel_kind<vector>,
+  pto.target_arch = "a5"
+} {
+  func.func @backend_attr_invalid() {
+    return
+  }
+}
+
+// CHECK: invalid pto.backend 'other'. Expected 'emitc' or 'vpto'.
diff --git a/test/lit/vpto/backend_attr_selects_vpto.pto b/test/lit/vpto/backend_attr_selects_vpto.pto
@@ -0,0 +1,26 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: ptoas --pto-arch=a5 --emit-vpto %s -o - | FileCheck %s
+
+module attributes {
+  pto.backend = "vpto",
+  pto.kernel_kind = #pto.kernel_kind<vector>,
+  pto.target_arch = "a5"
+} {
+  func.func @backend_attr_selects_vpto() {
+    %ctrl = pto.get_ctrl : i64
+    pto.set_ctrl %ctrl : i64
+    return
+  }
+}
+
+// CHECK: module
+// CHECK-SAME: pto.backend = "vpto"
+// CHECK: pto.get_ctrl
+// CHECK: pto.set_ctrl
diff --git a/test/lit/vpto/backend_child_attr_inherits_outer_arch.pto b/test/lit/vpto/backend_child_attr_inherits_outer_arch.pto
@@ -0,0 +1,29 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: ptoas --emit-vpto %s -o - | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func @backend_child_attr_inherits_outer_arch() {
+      %ctrl = pto.get_ctrl : i64
+      pto.set_ctrl %ctrl : i64
+      return
+    }
+  }
+}
+
+// CHECK: module
+// CHECK-SAME: pto.target_arch = "a5"
+// CHECK: module
+// CHECK-SAME: pto.backend = "vpto"
+// CHECK: pto.get_ctrl
+// CHECK: pto.set_ctrl
diff --git a/test/lit/vpto/backend_child_attr_selects_vpto.pto b/test/lit/vpto/backend_child_attr_selects_vpto.pto
@@ -0,0 +1,29 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: ptoas --pto-arch=a5 --emit-vpto %s -o - | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func @backend_child_attr_selects_vpto() {
+      %ctrl = pto.get_ctrl : i64
+      pto.set_ctrl %ctrl : i64
+      return
+    }
+  }
+}
+
+// CHECK: module
+// CHECK-SAME: pto.target_arch = "a5"
+// CHECK: module
+// CHECK-SAME: pto.backend = "vpto"
+// CHECK: pto.get_ctrl
+// CHECK: pto.set_ctrl
diff --git a/test/lit/vpto/backend_cli_override.pto b/test/lit/vpto/backend_cli_override.pto
@@ -0,0 +1,23 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: not ptoas --pto-arch=a5 --pto-backend=emitc --emit-vpto %s -o - 2>&1 | FileCheck %s
+
+module attributes {
+  pto.backend = "vpto",
+  pto.kernel_kind = #pto.kernel_kind<vector>,
+  pto.target_arch = "a5"
+} {
+  func.func @backend_cli_override() {
+    %ctrl = pto.get_ctrl : i64
+    pto.set_ctrl %ctrl : i64
+    return
+  }
+}
+
+// CHECK: Error: VPTO-specific flags require --pto-backend=vpto or pto.backend = "vpto".
diff --git a/test/lit/vpto/backend_cli_override_to_vpto.pto b/test/lit/vpto/backend_cli_override_to_vpto.pto
@@ -0,0 +1,24 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: ptoas --pto-arch=a5 --pto-backend=vpto --emit-vpto %s -o - | FileCheck %s
+
+module attributes {
+  pto.backend = "emitc",
+  pto.kernel_kind = #pto.kernel_kind<vector>,
+  pto.target_arch = "a5"
+} {
+  func.func @backend_cli_override_to_vpto() {
+    %ctrl = pto.get_ctrl : i64
+    pto.set_ctrl %ctrl : i64
+    return
+  }
+}
+
+// CHECK: pto.get_ctrl
+// CHECK: pto.set_ctrl
diff --git a/test/lit/vpto/backend_mixed_child_attrs_rejected.pto b/test/lit/vpto/backend_mixed_child_attrs_rejected.pto
@@ -0,0 +1,28 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: not ptoas --pto-arch=a5 --emit-vpto %s -o - 2>&1 | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module attributes {pto.backend = "emitc"} {
+    func.func @emitc_child() attributes {pto.aicore} {
+      return
+    }
+  }
+
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func @vpto_child() {
+      return
+    }
+  }
+}
+
+// CHECK: mixed pto.backend fatobj mode requires an explicit file path passed with -o
diff --git a/test/lit/vpto/backend_mixed_child_missing_backend_invalid.pto b/test/lit/vpto/backend_mixed_child_missing_backend_invalid.pto
@@ -0,0 +1,34 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: not ptoas --pto-arch=a5 --emit-vpto %s -o - 2>&1 | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module attributes {pto.backend = "emitc"} {
+    func.func @emitc_child() attributes {pto.aicore} {
+      return
+    }
+  }
+
+  module {
+    func.func @missing_backend_child() {
+      return
+    }
+  }
+
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func @vpto_child() {
+      return
+    }
+  }
+}
+
+// CHECK: mixed-backend container child module is missing pto.backend
diff --git a/test/lit/vpto/backend_mixed_duplicate_abi_export_invalid.pto b/test/lit/vpto/backend_mixed_duplicate_abi_export_invalid.pto
@@ -0,0 +1,37 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// 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.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: not ptoas --pto-arch=a5 %s -o %t.o 2>&1 | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module attributes {pto.backend = "emitc"} {
+    func.func @emitc_child() attributes {pto.aicore} {
+      return
+    }
+  }
+
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func public @vpto_post() {
+      return
+    }
+  }
+
+  module attributes {
+    pto.backend = "vpto",
+    pto.kernel_kind = #pto.kernel_kind<vector>
+  } {
+    func.func public @vpto_post() {
+      return
+    }
+  }
+}
+
+// CHECK: generated ABI export symbol 'vpto_post.vector' is defined by multiple backend child modules