From 47c51e8e38eee9c0179ca7c98d900515c7456cc1 Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 14 Apr 2026 20:02:37 +0800 Subject: [PATCH 1/2] fix: fall back illegal a5 auto sync pairs to barrier --- lib/PTO/Transforms/InsertSync/SyncCodegen.cpp | 42 +++++++++++++++++++ .../textract_a5_scaling_pipe_selection.pto | 5 ++- test/basic/tinsert_a5_pipe_selection.pto | 12 ++++-- test/basic/tmov_acc_mat_pipe_selection.pto | 14 +++++-- test/basic/tmov_acc_to_vec_mode_a5_emitc.pto | 5 ++- 5 files changed, 66 insertions(+), 12 deletions(-) diff --git a/lib/PTO/Transforms/InsertSync/SyncCodegen.cpp b/lib/PTO/Transforms/InsertSync/SyncCodegen.cpp index f6a4ff0c4..9ff917a5b 100644 --- a/lib/PTO/Transforms/InsertSync/SyncCodegen.cpp +++ b/lib/PTO/Transforms/InsertSync/SyncCodegen.cpp @@ -36,6 +36,26 @@ static pto::EventAttr getEventAttr(Builder &builder, int id) { auto odsEventVal = static_cast(id); return pto::EventAttr::get(builder.getContext(), odsEventVal); } + +static bool isA5LowLevelSyncPipeLegal(PipelineType pipe) { + switch (pipe) { + case PipelineType::PIPE_S: + case PipelineType::PIPE_V: + case PipelineType::PIPE_MTE2: + case PipelineType::PIPE_MTE3: + return true; + default: + return false; + } +} + +static bool shouldUseA5BarrierFallback(func::FuncOp func, + const SyncOperation *sync) { + if (!isTargetArchA5(func.getOperation())) + return false; + return !isA5LowLevelSyncPipeLegal(sync->GetActualSrcPipe()) || + !isA5LowLevelSyncPipeLegal(sync->GetActualDstPipe()); +} static bool IsSameSyncSignature(const SyncOperation *existing, const SyncOperation *candidate) { @@ -335,6 +355,17 @@ void SyncCodegen::CreateSetWaitOpForSingleBuffer(IRRewriter &rewriter, Operation *op, SyncOperation *sync, bool beforeInsert) { + if (shouldUseA5BarrierFallback(func_, sync)) { + auto pipeAllAttr = getPipeAttr(rewriter, PipelineType::PIPE_ALL); + if (beforeInsert || op->hasTrait()) { + rewriter.setInsertionPoint(op); + } else { + rewriter.setInsertionPointAfter(op); + } + rewriter.create(op->getLoc(), pipeAllAttr); + return; + } + // [Fix] Terminator 强制前置插入 if (beforeInsert || op->hasTrait()) { rewriter.setInsertionPoint(op); @@ -357,6 +388,17 @@ void SyncCodegen::CreateSetWaitOpForMultiBuffer(IRRewriter &rewriter, Operation *op, SyncOperation *sync, bool beforeInsert) { + if (shouldUseA5BarrierFallback(func_, sync)) { + auto pipeAllAttr = getPipeAttr(rewriter, PipelineType::PIPE_ALL); + if (beforeInsert || op->hasTrait()) { + rewriter.setInsertionPoint(op); + } else { + rewriter.setInsertionPointAfter(op); + } + rewriter.create(op->getLoc(), pipeAllAttr); + return; + } + // 注意：GetBufferSelected 可能需要在插入 Set/Wait 之前调用，以确保 SSA 顺序 // 但这里只是获取 Value，不影响 InsertionPoint 的设定 Value bufferSelected = GetBufferSelected(rewriter, op, sync); diff --git a/test/basic/textract_a5_scaling_pipe_selection.pto b/test/basic/textract_a5_scaling_pipe_selection.pto index 8a958fad5..8fa784071 100644 --- a/test/basic/textract_a5_scaling_pipe_selection.pto +++ b/test/basic/textract_a5_scaling_pipe_selection.pto @@ -17,7 +17,8 @@ module attributes {"pto.target_arch" = "a5"} { } // CHECK-LABEL: __global__ AICORE void textract_mat_scaling_sync( -// CHECK: set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0); -// CHECK: wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0); +// CHECK-NOT: set_flag(PIPE_MTE2, PIPE_MTE1 +// CHECK-NOT: wait_flag(PIPE_MTE2, PIPE_MTE1 // CHECK-NOT: set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); +// CHECK: pipe_barrier(PIPE_ALL); // CHECK: TEXTRACT( diff --git a/test/basic/tinsert_a5_pipe_selection.pto b/test/basic/tinsert_a5_pipe_selection.pto index a72d39d9b..950b2e55b 100644 --- a/test/basic/tinsert_a5_pipe_selection.pto +++ b/test/basic/tinsert_a5_pipe_selection.pto @@ -51,10 +51,14 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // CHECK-LABEL: __global__ AICORE void tinsert_acc_mat_pipeline( -// CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// CHECK-NOT: set_flag(PIPE_M, PIPE_FIX +// CHECK-NOT: wait_flag(PIPE_M, PIPE_FIX +// CHECK-NOT: set_flag(PIPE_FIX, PIPE_MTE1 +// CHECK-NOT: wait_flag(PIPE_FIX, PIPE_MTE1 +// CHECK-NOT: set_flag(PIPE_MTE1, PIPE_M +// CHECK-NOT: wait_flag(PIPE_MTE1, PIPE_M // CHECK-NOT: set_flag(PIPE_M, PIPE_MTE3, EVENT_ID0); -// CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// CHECK: pipe_barrier(PIPE_ALL); // CHECK: TINSERT( -// CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); +// CHECK: pipe_barrier(PIPE_ALL); // CHECK-NOT: set_flag(PIPE_MTE3, PIPE_MTE1, EVENT_ID0); -// CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); diff --git a/test/basic/tmov_acc_mat_pipe_selection.pto b/test/basic/tmov_acc_mat_pipe_selection.pto index cba124563..7a5f1e3ce 100644 --- a/test/basic/tmov_acc_mat_pipe_selection.pto +++ b/test/basic/tmov_acc_mat_pipe_selection.pto @@ -32,8 +32,14 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // CHECK-LABEL: __global__ AICORE void tmov_acc_mat_pipeline( -// CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); -// CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// CHECK-NOT: set_flag(PIPE_M, PIPE_FIX +// CHECK-NOT: wait_flag(PIPE_M, PIPE_FIX +// CHECK-NOT: set_flag(PIPE_FIX, PIPE_MTE1 +// CHECK-NOT: wait_flag(PIPE_FIX, PIPE_MTE1 +// CHECK-NOT: set_flag(PIPE_MTE1, PIPE_M +// CHECK-NOT: wait_flag(PIPE_MTE1, PIPE_M +// CHECK: pipe_barrier(PIPE_ALL); // CHECK: TMOV( -// CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); -// CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); +// CHECK: pipe_barrier(PIPE_ALL); +// CHECK: TMATMUL( +// CHECK: pipe_barrier(PIPE_ALL); diff --git a/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto b/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto index 8800e9af3..234f290e6 100644 --- a/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto +++ b/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto @@ -30,8 +30,9 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // A5-LABEL: __global__ AICORE void tmov_acc_to_vec_mode_a5( -// A5: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); -// A5: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// A5-NOT: set_flag(PIPE_M, PIPE_FIX +// A5-NOT: wait_flag(PIPE_M, PIPE_FIX +// A5: pipe_barrier(PIPE_ALL); // A5: TMOV<{{.*}}pto::AccToVecMode::DualModeSplitM{{.*}}>( // A3-LABEL: __global__ AICORE void tmov_acc_to_vec_mode_a5( // A3: TMOV<{{.*}}pto::AccToVecMode::DualModeSplitM{{.*}}>( From 7b45c747a15f306543fd65a62e9d33d8122b6eef Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 14 Apr 2026 20:19:54 +0800 Subject: [PATCH 2/2] chore: drop basic case updates from a5 sync fix --- test/basic/textract_a5_scaling_pipe_selection.pto | 5 ++--- test/basic/tinsert_a5_pipe_selection.pto | 12 ++++-------- test/basic/tmov_acc_mat_pipe_selection.pto | 14 ++++---------- test/basic/tmov_acc_to_vec_mode_a5_emitc.pto | 5 ++--- 4 files changed, 12 insertions(+), 24 deletions(-) diff --git a/test/basic/textract_a5_scaling_pipe_selection.pto b/test/basic/textract_a5_scaling_pipe_selection.pto index 8fa784071..8a958fad5 100644 --- a/test/basic/textract_a5_scaling_pipe_selection.pto +++ b/test/basic/textract_a5_scaling_pipe_selection.pto @@ -17,8 +17,7 @@ module attributes {"pto.target_arch" = "a5"} { } // CHECK-LABEL: __global__ AICORE void textract_mat_scaling_sync( -// CHECK-NOT: set_flag(PIPE_MTE2, PIPE_MTE1 -// CHECK-NOT: wait_flag(PIPE_MTE2, PIPE_MTE1 +// CHECK: set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0); +// CHECK: wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0); // CHECK-NOT: set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); -// CHECK: pipe_barrier(PIPE_ALL); // CHECK: TEXTRACT( diff --git a/test/basic/tinsert_a5_pipe_selection.pto b/test/basic/tinsert_a5_pipe_selection.pto index 950b2e55b..a72d39d9b 100644 --- a/test/basic/tinsert_a5_pipe_selection.pto +++ b/test/basic/tinsert_a5_pipe_selection.pto @@ -51,14 +51,10 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // CHECK-LABEL: __global__ AICORE void tinsert_acc_mat_pipeline( -// CHECK-NOT: set_flag(PIPE_M, PIPE_FIX -// CHECK-NOT: wait_flag(PIPE_M, PIPE_FIX -// CHECK-NOT: set_flag(PIPE_FIX, PIPE_MTE1 -// CHECK-NOT: wait_flag(PIPE_FIX, PIPE_MTE1 -// CHECK-NOT: set_flag(PIPE_MTE1, PIPE_M -// CHECK-NOT: wait_flag(PIPE_MTE1, PIPE_M +// CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); // CHECK-NOT: set_flag(PIPE_M, PIPE_MTE3, EVENT_ID0); -// CHECK: pipe_barrier(PIPE_ALL); +// CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); // CHECK: TINSERT( -// CHECK: pipe_barrier(PIPE_ALL); +// CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); // CHECK-NOT: set_flag(PIPE_MTE3, PIPE_MTE1, EVENT_ID0); +// CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); diff --git a/test/basic/tmov_acc_mat_pipe_selection.pto b/test/basic/tmov_acc_mat_pipe_selection.pto index 7a5f1e3ce..cba124563 100644 --- a/test/basic/tmov_acc_mat_pipe_selection.pto +++ b/test/basic/tmov_acc_mat_pipe_selection.pto @@ -32,14 +32,8 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // CHECK-LABEL: __global__ AICORE void tmov_acc_mat_pipeline( -// CHECK-NOT: set_flag(PIPE_M, PIPE_FIX -// CHECK-NOT: wait_flag(PIPE_M, PIPE_FIX -// CHECK-NOT: set_flag(PIPE_FIX, PIPE_MTE1 -// CHECK-NOT: wait_flag(PIPE_FIX, PIPE_MTE1 -// CHECK-NOT: set_flag(PIPE_MTE1, PIPE_M -// CHECK-NOT: wait_flag(PIPE_MTE1, PIPE_M -// CHECK: pipe_barrier(PIPE_ALL); +// CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); // CHECK: TMOV( -// CHECK: pipe_barrier(PIPE_ALL); -// CHECK: TMATMUL( -// CHECK: pipe_barrier(PIPE_ALL); +// CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); +// CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0); diff --git a/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto b/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto index 234f290e6..8800e9af3 100644 --- a/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto +++ b/test/basic/tmov_acc_to_vec_mode_a5_emitc.pto @@ -30,9 +30,8 @@ module attributes {"pto.device-spec" = "Ascend950"} { } // A5-LABEL: __global__ AICORE void tmov_acc_to_vec_mode_a5( -// A5-NOT: set_flag(PIPE_M, PIPE_FIX -// A5-NOT: wait_flag(PIPE_M, PIPE_FIX -// A5: pipe_barrier(PIPE_ALL); +// A5: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0); +// A5: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0); // A5: TMOV<{{.*}}pto::AccToVecMode::DualModeSplitM{{.*}}>( // A3-LABEL: __global__ AICORE void tmov_acc_to_vec_mode_a5( // A3: TMOV<{{.*}}pto::AccToVecMode::DualModeSplitM{{.*}}>(