mouliangyu · mouliangyu · May 29, 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
@@ -180,12 +180,20 @@ class ScopedPTOParserTargetArch {
 /// Function attribute that marks an explicit PTO kernel entry.
 inline constexpr llvm::StringLiteral kPTOEntryAttrName = "pto.entry";
 inline constexpr llvm::StringLiteral kLegacyHACCEntryAttrName = "hacc.entry";
+inline constexpr llvm::StringLiteral kPTOKernelAttrName = "pto.kernel";
+inline constexpr llvm::StringLiteral kLegacyPTOAICoreAttrName = "pto.aicore";
 inline constexpr llvm::StringLiteral kPTOSimtEntryAttrName = "pto.simt_entry";
 inline constexpr llvm::StringLiteral kPTOSimtMaxThreadsAttrName =
     "pto.simt_max_threads";
 inline constexpr llvm::StringLiteral kPTOSimtMaxRegistersAttrName =
     "pto.simt_max_regs";
 
+/// Return true if the operation carries a PTO kernel marker.
+bool hasPTOKernelAttr(Operation *op);
+
+/// Return true if the function is a PTO kernel definition.
+bool isPTOKernelFunction(func::FuncOp func);
+
 /// Return true if the function carries an explicit entry marker.
 bool hasExplicitPTOEntryAttr(func::FuncOp func);
 

diff --git a/include/PTO/Transforms/VPTOLLVMEmitter.h b/include/PTO/Transforms/VPTOLLVMEmitter.h
@@ -37,6 +37,11 @@ struct VPTOEmissionOptions {
 };
 
 struct EmittedLLVMModule {
+  void reset() {
+    module.reset();
+    context.reset();
+  }
+
   std::unique_ptr<llvm::LLVMContext> context;
   std::unique_ptr<llvm::Module> module;
 };

diff --git a/lib/PTO/IR/PTO.cpp b/lib/PTO/IR/PTO.cpp
@@ -2125,6 +2125,15 @@ bool mlir::pto::isScalarPtrOrMemRef(Type type) {
   return false;
 }
 
+bool mlir::pto::hasPTOKernelAttr(Operation *op) {
+  return op && (op->hasAttr(kPTOKernelAttrName) ||
+                op->hasAttr(kLegacyPTOAICoreAttrName));
+}
+
+bool mlir::pto::isPTOKernelFunction(func::FuncOp func) {
+  return func && !func.isDeclaration() && hasPTOKernelAttr(func.getOperation());
+}
+
 bool mlir::pto::hasExplicitPTOEntryAttr(func::FuncOp func) {
   return func && (func->hasAttrOfType<UnitAttr>(kPTOEntryAttrName) ||
                   func->hasAttrOfType<UnitAttr>(kLegacyHACCEntryAttrName));

diff --git a/lib/PTO/Transforms/PTOToEmitC.cpp b/lib/PTO/Transforms/PTOToEmitC.cpp
@@ -2740,14 +2740,15 @@ struct FuncToEmitC : public OpConversionPattern<func::FuncOp> {
     }
 
     if (op.isDeclaration()) {
-      emitcFunc.setSpecifiersAttr(rewriter.getStrArrayAttr({"extern"}));
+      emitcFunc.setSpecifiersAttr(
+          rewriter.getStrArrayAttr({"extern \"C\"", "AICORE"}));
       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/PTOViewToMemref.cpp b/lib/PTO/Transforms/PTOViewToMemref.cpp
@@ -1174,9 +1174,6 @@ struct PTOViewToMemrefPass
     MLIRContext *ctx = &getContext();
 
     for (auto func : mod.getOps<func::FuncOp>()) {
-      if (func.isExternal()) continue;
-
-      Block &entry = func.front();
       auto fnTy = func.getFunctionType();
 
       // ------------------------------------------------------------------
@@ -1188,16 +1185,18 @@ struct PTOViewToMemrefPass
       SmallVector<Type> newResults;
       for (Type t : fnTy.getResults()) newResults.push_back(convertPTOTypeToMemRef(t));
 
+      func.setFunctionType(FunctionType::get(ctx, newInputs, newResults));
+      if (func.isExternal()) continue;
+
+      Block &entry = func.front();
+
       // Update entry block arguments
       for (unsigned i = 0; i < entry.getNumArguments(); ++i) {
         if (entry.getArgument(i).getType() != newInputs[i]) {
             entry.getArgument(i).setType(newInputs[i]);
         }
       }
 
-      // Update function type
-      func.setFunctionType(FunctionType::get(ctx, newInputs, newResults));
-
       // ------------------------------------------------------------------
       // Stage 0.25: Insert pto.bind_tile for function args that were tile_buf
       // ------------------------------------------------------------------

diff --git a/lib/PTO/Transforms/VPTOLLVMEmitter.cpp b/lib/PTO/Transforms/VPTOLLVMEmitter.cpp
@@ -59,16 +59,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);
@@ -9668,7 +9661,7 @@ static LogicalResult renameKernelFunctionsForKernelKind(ModuleOp module,
   }
 
   for (func::FuncOp funcOp : module.getOps<func::FuncOp>()) {
-    if (!hasVPTOKernelAttr(funcOp))
+    if (!pto::hasPTOKernelAttr(funcOp.getOperation()))
       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 does not support debug IR output flags
diff --git a/test/lit/vpto/backend_mixed_child_missing_backend_defaults.pto b/test/lit/vpto/backend_mixed_child_missing_backend_defaults.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 %s -o - 2>&1 | FileCheck %s
+
+module attributes {pto.target_arch = "a5"} {
+  module {
+    func.func @missing_backend_child_defaults_to_emitc() 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