fix(a2a3): give the V2C ring its own base in a bidirectional TPipe-pto-isa-AtomGit

已合并

fix(a2a3): give the V2C ring its own base in a bidirectional TPipe #1438

puff_the_magic_dragon创建于 16 天前从puff_the_magic_dragon/pto-isa： fix/a2a3-dir-both-ring-offset 合入到 cann/pto-isa： master

fix(a2a3): give the V2C ring its own base in a bidirectional TPipe #1438

已合并

puff_the_magic_dragon创建于 16 天前

从puff_the_magic_dragon/pto-isa fix/a2a3-dir-both-ring-offset合入到cann/pto-isa master

共 7 个文件变更+625-2

Minclude/pto/npu/a2a3/TPush.hpp+16-1

	@@ -469,7 +469,22 @@ struct TPipe {
469		469
470	PTO_INTERNAL explicit TPipe(__gm__ void* GM_SLOT_BUFFER, uint32_t C2V_CONSUMER_BUF, uint32_t V2C_CONSUMER_BUF)	470	PTO_INTERNAL explicit TPipe(__gm__ void* GM_SLOT_BUFFER, uint32_t C2V_CONSUMER_BUF, uint32_t V2C_CONSUMER_BUF)
471	: fifo(GM_SLOT_BUFFER, C2V_CONSUMER_BUF, V2C_CONSUMER_BUF), prod(), cons()	471	: fifo(GM_SLOT_BUFFER, C2V_CONSUMER_BUF, V2C_CONSUMER_BUF), prod(), cons()
472	- {}	472	+ {
		473	+ // Bidirectional: the two rings must NOT overlap. Both directions index the shared GM
		474	+ // buffer as (tileIndex % SlotNum) * SlotSize, so without a per-direction base the C2V
		475	+ // and V2C rings alias the same slots and silently corrupt each other's tiles. Place
		476	+ // V2C after C2V, exactly as the ISA reference specifies
		477	+ // (`v2c_ring_buf = GM_SLOT_BUFFER + SLOT_NUM * SLOT_SIZE`).
		478	+ if constexpr (is_both) {
		479	+ constexpr int V2C_ENTRY_OFFSET = static_cast<int>(SlotNum) * static_cast<int>(SlotSize);
		480	+#ifdef __DAV_CUBE__
		481	+ cons.setEntryOffset(V2C_ENTRY_OFFSET); // Cube consumes V2C
		482	+#endif
		483	+#ifdef __DAV_VEC__
		484	+ prod.setEntryOffset(V2C_ENTRY_OFFSET); // Vector produces V2C
		485	+#endif
		486	+ }
		487	+ }
473		488
474	// Destructor for TPipe: drain leftover free credits on FlagID+1.	489	// Destructor for TPipe: drain leftover free credits on FlagID+1.
475	// Initial TPUSH calls skip allocate() via shouldWaitFree (tileIndex < SlotNum, or 0 for depth 1).	490	// Initial TPUSH calls skip allocate() via shouldWaitFree (tileIndex < SlotNum, or 0 for depth 1).

Mtests/npu/a2a3/src/st/testcase/CMakeLists.txt+2-0

	@@ -226,6 +226,7 @@ tpushpop_vc
226	tpushpop_cv_nosplit	226	tpushpop_cv_nosplit
227	tpushpop_vc_nosplit	227	tpushpop_vc_nosplit
228	tpushpop_dir_both	228	tpushpop_dir_both
		229	+tpushpop_dir_both_concurrent
229	tpushpop_subtile	230	tpushpop_subtile
230	tfusedmuladd	231	tfusedmuladd
231	tmuladddst	232	tmuladddst
	@@ -254,6 +255,7 @@ if (AUTO_MODE)
254	tpushpop_cv_nosplit	255	tpushpop_cv_nosplit
255	tpushpop_vc_nosplit	256	tpushpop_vc_nosplit
256	tpushpop_dir_both	257	tpushpop_dir_both
		258	+ tpushpop_dir_both_concurrent
257	tpushpop_subtile	259	tpushpop_subtile
258	# the following case wait to be fix	260	# the following case wait to be fix
259	mgather	261	mgather

Mtests/npu/a2a3/src/st/testcase/tpushpop_dir_both/main.cpp+4-1

	@@ -44,7 +44,10 @@ void TPushPopDirBothTestFunc(uint32_t M, uint32_t K, uint32_t N)
44	size_t dFileSize = K * N * sizeof(T);	44	size_t dFileSize = K * N * sizeof(T);
45	size_t fFileSize = M * N * sizeof(T);	45	size_t fFileSize = M * N * sizeof(T);
46	size_t outFileSize = M * N * sizeof(T);	46	size_t outFileSize = M * N * sizeof(T);
47	- size_t fifoFileSize = 2 * M * N * sizeof(T);	47	+ // A DIR_BOTH pipe is two rings in GM: C2V at offset 0 and V2C at SLOT_NUM * SLOT_SIZE.
		48	+ // SLOT_SIZE = M * N * sizeof(T) and SLOT_NUM = FIFO_DEPTH = 2 in the kernel, so the
		49	+ // buffer must be 2 * SLOT_NUM * SLOT_SIZE.
		50	+ size_t fifoFileSize = 4 * M * N * sizeof(T);
48		51
49	aclInit(nullptr);	52	aclInit(nullptr);
50	aclrtSetDevice(0);	53	aclrtSetDevice(0);

Atests/npu/a2a3/src/st/testcase/tpushpop_dir_both_concurrent/CMakeLists.txt+11-0

@@ -0,0 +1,11 @@
	1	+# --------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# --------------------------------------------------------------------------------
	10	+
	11	+pto_mix_st(tpushpop_dir_both_concurrent)

Atests/npu/a2a3/src/st/testcase/tpushpop_dir_both_concurrent/gen_data.py+72-0

@@ -0,0 +1,72 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+# --------------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# --------------------------------------------------------------------------------
	12	+
	13	+import os
	14	+import numpy as np
	15	+
	16	+np.random.seed(19)
	17	+
	18	+
	19	+def gen_golden_data(case_name, case_params):
	20	+ """
	21	+ DIR_BOTH concurrent test: the two directions are IN FLIGHT AT THE SAME TIME.
	22	+
	23	+ Unlike tpushpop_dir_both, the cube's matmul takes both operands from GM, so the
	24	+ C2V push does not depend on the V2C pop and is issued first. The vector still
	25	+ pushes V2C then pops C2V. Both tiles are tileIndex 0 -> both address slot 0.
	26	+
	27	+ Computation:
	28	+ C = A + B [M, K] (vector -> cube, V2C; never consumed here)
	29	+ E = A @ D [M, N] (cube -> vector, C2V; GM operands only)
	30	+ G = E - F [M, N] (golden output)
	31	+ """
	32	+ m, k, n, dtype = case_params
	33	+ srcA = np.random.uniform(-2, 2, [m, k]).astype(dtype)
	34	+ srcB = np.random.uniform(-2, 2, [m, k]).astype(dtype)
	35	+ srcD = np.random.uniform(-2, 2, [k, n]).astype(dtype)
	36	+ srcF = np.random.uniform(-2, 2, [m, n]).astype(dtype)
	37	+
	38	+ srcA.tofile("./srcA_gm.bin")
	39	+ srcB.tofile("./srcB_gm.bin")
	40	+ srcD.tofile("./srcD_gm.bin")
	41	+ srcF.tofile("./srcF_gm.bin")
	42	+
	43	+ e_mat = np.matmul(srcA, srcD).astype(dtype)
	44	+ golden = (e_mat - srcF).astype(dtype)
	45	+ golden.tofile("./golden.bin")
	46	+
	47	+ # What the CUBE must receive over V2C, put through the same second matmul the kernel
	48	+ # applies to it. Checking this alongside the vector's output is what makes the slot-0
	49	+ # aliasing observable -- a single-sided check can still pass on an affected build.
	50	+ c_mat = (srcA + srcB).astype(dtype)
	51	+ golden_cube = np.matmul(c_mat, srcD).astype(dtype)
	52	+ golden_cube.tofile("./golden_cube.bin")
	53	+
	54	+
	55	+if __name__ == "__main__":
	56	+ case_name_list = [
	57	+ "TPushPopDirBothConcurrentTest.case1_float_dir_both_concurrent",
	58	+ "TPushPopDirBothConcurrentTest.case2_float_dir_both_concurrent_left_right",
	59	+ ]
	60	+
	61	+ case_params_list = [
	62	+ (128, 64, 128, np.float32),
	63	+ (128, 64, 128, np.float32),
	64	+ ]
	65	+
	66	+ for i, case_name in enumerate(case_name_list):
	67	+ if not os.path.exists(case_name):
	68	+ os.makedirs(case_name)
	69	+ original_dir = os.getcwd()
	70	+ os.chdir(case_name)
	71	+ gen_golden_data(case_name, case_params_list[i])
	72	+ os.chdir(original_dir)

Atests/npu/a2a3/src/st/testcase/tpushpop_dir_both_concurrent/main.cpp+170-0

@@ -0,0 +1,170 @@
	1	+/**
	2	+Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+CANN Open Software License Agreement Version 2.0 (the "License").
	5	+Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+See LICENSE in the root of the software repository for the full text of the License.
	9	+*/
	10	+
	11	+#include "test_common.h"
	12	+#include "acl/acl.h"
	13	+#include "runtime/rt.h"
	14	+#include <gtest/gtest.h>
	15	+
	16	+using namespace std;
	17	+using namespace PtoTestCommon;
	18	+
	19	+template <int32_t tilingKey>
	20	+void LaunchTPushPopDirBoth(
	21	+ uint8_t* ffts, uint8_t* out, uint8_t* srcA, uint8_t* srcB, uint8_t* srcD, uint8_t* srcF, uint8_t* fifoMem,
	22	+ uint8_t* outCube, void* stream);
	23	+
	24	+class TPushPopDirBothConcurrentTest : public testing::Test {
	25	+protected:
	26	+ void SetUp() override {}
	27	+ void TearDown() override {}
	28	+};
	29	+
	30	+std::string GetGoldenDir()
	31	+{
	32	+ const testing::TestInfo* testInfo = testing::UnitTest::GetInstance()->current_test_info();
	33	+ const std::string caseName = testInfo->name();
	34	+ std::string suiteName = testInfo->test_suite_name();
	35	+ std::string fullPath = "../" + suiteName + "." + caseName;
	36	+ return fullPath;
	37	+}
	38	+
	39	+// Host and device buffers for one run, grouped so setup / teardown / verify can be separate
	40	+// functions instead of one long test body.
	41	+struct DirBothBuffers {
	42	+ size_t aSize, bSize, dSize, fSize, outSize, fifoSize;
	43	+ uint8_t outHost, outCubeHost, srcAHost, srcBHost, srcDHost, srcFHost;
	44	+ uint8_t outDevice, outCubeDevice, srcADevice, srcBDevice, srcDDevice, srcFDevice, *fifoMemDevice;
	45	+};
	46	+
	47	+template <typename T>
	48	+void SetupBuffers(DirBothBuffers& b, uint32_t M, uint32_t K, uint32_t N)
	49	+{
	50	+ b.aSize = M * K * sizeof(T);
	51	+ b.bSize = M * K * sizeof(T);
	52	+ b.dSize = K * N * sizeof(T);
	53	+ b.fSize = M * N * sizeof(T);
	54	+ b.outSize = M * N * sizeof(T);
	55	+ // A DIR_BOTH pipe is two rings in GM: C2V at offset 0 and V2C at SLOT_NUM * SLOT_SIZE.
	56	+ // SLOT_SIZE = M * N * sizeof(T) and SLOT_NUM = FIFO_DEPTH = 2 in the kernel, so the
	57	+ // buffer must be 2 * SLOT_NUM * SLOT_SIZE.
	58	+ b.fifoSize = 4 * M * N * sizeof(T);
	59	+
	60	+ aclrtMallocHost((void**)(&b.outHost), b.outSize);
	61	+ aclrtMallocHost((void**)(&b.outCubeHost), b.outSize);
	62	+ aclrtMallocHost((void**)(&b.srcAHost), b.aSize);
	63	+ aclrtMallocHost((void**)(&b.srcBHost), b.bSize);
	64	+ aclrtMallocHost((void**)(&b.srcDHost), b.dSize);
	65	+ aclrtMallocHost((void**)(&b.srcFHost), b.fSize);
	66	+
	67	+ aclrtMalloc((void**)&b.outDevice, b.outSize, ACL_MEM_MALLOC_HUGE_FIRST);
	68	+ aclrtMalloc((void**)&b.outCubeDevice, b.outSize, ACL_MEM_MALLOC_HUGE_FIRST);
	69	+ aclrtMalloc((void**)&b.srcADevice, b.aSize, ACL_MEM_MALLOC_HUGE_FIRST);
	70	+ aclrtMalloc((void**)&b.srcBDevice, b.bSize, ACL_MEM_MALLOC_HUGE_FIRST);
	71	+ aclrtMalloc((void**)&b.srcDDevice, b.dSize, ACL_MEM_MALLOC_HUGE_FIRST);
	72	+ aclrtMalloc((void**)&b.srcFDevice, b.fSize, ACL_MEM_MALLOC_HUGE_FIRST);
	73	+ aclrtMalloc((void**)&b.fifoMemDevice, b.fifoSize, ACL_MEM_MALLOC_HUGE_FIRST);
	74	+
	75	+ ReadFile(GetGoldenDir() + "/srcA_gm.bin", b.aSize, b.srcAHost, b.aSize);
	76	+ ReadFile(GetGoldenDir() + "/srcB_gm.bin", b.bSize, b.srcBHost, b.bSize);
	77	+ ReadFile(GetGoldenDir() + "/srcD_gm.bin", b.dSize, b.srcDHost, b.dSize);
	78	+ ReadFile(GetGoldenDir() + "/srcF_gm.bin", b.fSize, b.srcFHost, b.fSize);
	79	+
	80	+ aclrtMemcpy(b.srcADevice, b.aSize, b.srcAHost, b.aSize, ACL_MEMCPY_HOST_TO_DEVICE);
	81	+ aclrtMemcpy(b.srcBDevice, b.bSize, b.srcBHost, b.bSize, ACL_MEMCPY_HOST_TO_DEVICE);
	82	+ aclrtMemcpy(b.srcDDevice, b.dSize, b.srcDHost, b.dSize, ACL_MEMCPY_HOST_TO_DEVICE);
	83	+ aclrtMemcpy(b.srcFDevice, b.fSize, b.srcFHost, b.fSize, ACL_MEMCPY_HOST_TO_DEVICE);
	84	+}
	85	+
	86	+void TeardownBuffers(DirBothBuffers& b)
	87	+{
	88	+ aclrtFree(b.outDevice);
	89	+ aclrtFree(b.outCubeDevice);
	90	+ aclrtFree(b.srcADevice);
	91	+ aclrtFree(b.srcBDevice);
	92	+ aclrtFree(b.srcDDevice);
	93	+ aclrtFree(b.srcFDevice);
	94	+ aclrtFree(b.fifoMemDevice);
	95	+
	96	+ aclrtFreeHost(b.outHost);
	97	+ aclrtFreeHost(b.outCubeHost);
	98	+ aclrtFreeHost(b.srcAHost);
	99	+ aclrtFreeHost(b.srcBHost);
	100	+ aclrtFreeHost(b.srcDHost);
	101	+ aclrtFreeHost(b.srcFHost);
	102	+}
	103	+
	104	+// Check BOTH directions. The two payloads are different sizes, so whichever push lands
	105	+// second overwrites only part of the other tile and a single-sided check can still come
	106	+// back clean. On an affected build both comparisons fail.
	107	+template <typename T>
	108	+void VerifyOutputs(size_t outFileSize)
	109	+{
	110	+ std::vector<T> golden(outFileSize);
	111	+ std::vector<T> devFinal(outFileSize);
	112	+ ReadFile(GetGoldenDir() + "/golden.bin", outFileSize, golden.data(), outFileSize);
	113	+ ReadFile(GetGoldenDir() + "/output_z.bin", outFileSize, devFinal.data(), outFileSize);
	114	+
	115	+ bool ret = ResultCmp(golden, devFinal, 0.001f);
	116	+ EXPECT_TRUE(ret) << "vector side (C2V) mismatch";
	117	+
	118	+ std::vector<T> goldenCube(outFileSize);
	119	+ std::vector<T> devCube(outFileSize);
	120	+ ReadFile(GetGoldenDir() + "/golden_cube.bin", outFileSize, goldenCube.data(), outFileSize);
	121	+ ReadFile(GetGoldenDir() + "/output_cube.bin", outFileSize, devCube.data(), outFileSize);
	122	+
	123	+ bool retCube = ResultCmp(goldenCube, devCube, 0.001f);
	124	+ EXPECT_TRUE(retCube) << "cube side (V2C) mismatch -- C2V and V2C aliased the same GM slot";
	125	+}
	126	+
	127	+template <typename T, int32_t key>
	128	+void TPushPopDirBothConcurrentTestFunc(uint32_t M, uint32_t K, uint32_t N)
	129	+{
	130	+ aclInit(nullptr);
	131	+ aclrtSetDevice(0);
	132	+ aclrtStream stream;
	133	+ aclrtCreateStream(&stream);
	134	+
	135	+ DirBothBuffers b{};
	136	+ SetupBuffers<T>(b, M, K, N);
	137	+
	138	+ uint64_t ffts{0};
	139	+ uint32_t fftsLen{0};
	140	+ rtGetC2cCtrlAddr(&ffts, &fftsLen);
	141	+
	142	+ LaunchTPushPopDirBoth<key>(
	143	+ (uint8_t*)ffts, b.outDevice, b.srcADevice, b.srcBDevice, b.srcDDevice, b.srcFDevice, b.fifoMemDevice,
	144	+ b.outCubeDevice, stream);
	145	+
	146	+ aclrtSynchronizeStream(stream);
	147	+ aclrtMemcpy(b.outHost, b.outSize, b.outDevice, b.outSize, ACL_MEMCPY_DEVICE_TO_HOST);
	148	+ aclrtMemcpy(b.outCubeHost, b.outSize, b.outCubeDevice, b.outSize, ACL_MEMCPY_DEVICE_TO_HOST);
	149	+
	150	+ WriteFile(GetGoldenDir() + "/output_z.bin", b.outHost, b.outSize);
	151	+ WriteFile(GetGoldenDir() + "/output_cube.bin", b.outCubeHost, b.outSize);
	152	+
	153	+ const size_t outFileSize = b.outSize;
	154	+ TeardownBuffers(b);
	155	+ aclrtDestroyStream(stream);
	156	+ aclrtResetDevice(0);
	157	+ aclFinalize();
	158	+
	159	+ VerifyOutputs<T>(outFileSize);
	160	+}
	161	+
	162	+TEST_F(TPushPopDirBothConcurrentTest, case1_float_dir_both_concurrent)
	163	+{
	164	+ TPushPopDirBothConcurrentTestFunc<float, 1>(128, 64, 128);
	165	+}
	166	+
	167	+TEST_F(TPushPopDirBothConcurrentTest, case2_float_dir_both_concurrent_left_right)
	168	+{
	169	+ TPushPopDirBothConcurrentTestFunc<float, 2>(128, 64, 128);
	170	+}

Atests/npu/a2a3/src/st/testcase/tpushpop_dir_both_concurrent/tpushpop_dir_both_concurrent_kernel.cpp+350-0

@@ -0,0 +1,350 @@
	1	+/**
	2	+Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+CANN Open Software License Agreement Version 2.0 (the "License").
	5	+Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+See LICENSE in the root of the software repository for the full text of the License.
	9	+*/
	10	+
	11	+#include <pto/pto-inst.hpp>
	12	+#include <pto/common/fifo.hpp>
	13	+
	14	+using namespace pto;
	15	+
	16	+#define VEC_CORES 2
	17	+
	18	+#ifdef __DAV_CUBE__
	19	+constexpr bool DAV_CUBE = true;
	20	+#else
	21	+constexpr bool DAV_CUBE = false;
	22	+#endif
	23	+
	24	+#ifdef __DAV_VEC__
	25	+constexpr bool DAV_VEC = true;
	26	+#else
	27	+constexpr bool DAV_VEC = false;
	28	+#endif
	29	+
	30	+template <typename T>
	31	+AICORE constexpr inline T CeilAlign(T num_1, T num_2)
	32	+{
	33	+ if (num_2 == 0) {
	34	+ return 0;
	35	+ }
	36	+ return (num_1 + num_2 - 1) / num_2 * num_2;
	37	+}
	38	+
	39	+// ---------------------------------------------------------------------------------------
	40	+// Vector side, phase 1: build tileC = tileA + tileB and push it V2C.
	41	+//
	42	+// FLAG CONTRACT: on return, (PIPE_V, PIPE_MTE2, EVENT_ID1) and (PIPE_MTE3, PIPE_V,
	43	+// EVENT_ID1) are left ARMED for runVecPopC2V, which consumes them. The two halves are one
	44	+// flag ledger split across two functions; keep them balanced as a pair.
	45	+// ---------------------------------------------------------------------------------------
	46	+template <typename T, int TOTAL_M, int K, int N, TileSplitAxis SplitAxis, typename PipeT>
	47	+AICORE inline void runVecPushV2C(PipeT& pipe, __gm__ T* srcA, __gm__ T* srcB, uint32_t subBlockIdx)
	48	+{
	49	+ constexpr uint32_t V2C_ROWS = (SplitAxis == TileSplitAxis::TILE_UP_DOWN) ? (TOTAL_M / VEC_CORES) : TOTAL_M;
	50	+ constexpr uint32_t V2C_COLS = (SplitAxis == TileSplitAxis::TILE_LEFT_RIGHT) ? (K / VEC_CORES) : K;
	51	+
	52	+ using VecTileK = Tile<TileType::Vec, T, V2C_ROWS, V2C_COLS, BLayout::RowMajor, V2C_ROWS, V2C_COLS>;
	53	+ using GlobalAB = GlobalTensor<
	54	+ T, pto::Shape<1, 1, 1, V2C_ROWS, V2C_COLS>, pto::Stride<TOTAL_M * K, TOTAL_M * K, V2C_ROWS * V2C_COLS, K, 1>>;
	55	+
	56	+ VecTileK tileA, tileB, tileC;
	57	+ TASSIGN(tileA, 0x0);
	58	+ TASSIGN(tileB, 0x4000);
	59	+ TASSIGN(tileC, 0x8000);
	60	+
	61	+ size_t abOffset;
	62	+ if constexpr (SplitAxis == TileSplitAxis::TILE_UP_DOWN) {
	63	+ abOffset = static_cast<size_t>(subBlockIdx * V2C_ROWS) * K;
	64	+ } else {
	65	+ abOffset = static_cast<size_t>(subBlockIdx) * V2C_COLS;
	66	+ }
	67	+
	68	+ GlobalAB globalA(srcA + abOffset);
	69	+ GlobalAB globalB(srcB + abOffset);
	70	+
	71	+ set_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	72	+ set_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	73	+
	74	+ wait_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	75	+
	76	+ TLOAD(tileA, globalA);
	77	+ TLOAD(tileB, globalB);
	78	+
	79	+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
	80	+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
	81	+
	82	+ wait_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	83	+
	84	+ TADD(tileC, tileA, tileB);
	85	+
	86	+ set_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	87	+
	88	+ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
	89	+ wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
	90	+
	91	+ TPUSH<PipeT, VecTileK, SplitAxis>(pipe, tileC);
	92	+
	93	+ set_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	94	+}
	95	+
	96	+// ---------------------------------------------------------------------------------------
	97	+// Vector side, phase 2: pop the C2V tile, subtract tileF, store. Consumes the two flags
	98	+// runVecPushV2C left armed (see its FLAG CONTRACT) and drains the ledger before returning.
	99	+// ---------------------------------------------------------------------------------------
	100	+template <typename T, int TOTAL_M, int K, int N, TileSplitAxis SplitAxis, typename PipeT>
	101	+AICORE inline void runVecPopC2V(PipeT& pipe, __gm__ T* out, __gm__ T* srcF, uint32_t subBlockIdx)
	102	+{
	103	+ constexpr uint32_t C2V_ROWS = (SplitAxis == TileSplitAxis::TILE_UP_DOWN) ? (TOTAL_M / VEC_CORES) : TOTAL_M;
	104	+ constexpr uint32_t C2V_COLS = (SplitAxis == TileSplitAxis::TILE_LEFT_RIGHT) ? (N / VEC_CORES) : N;
	105	+
	106	+ using VecTileN = Tile<TileType::Vec, T, C2V_ROWS, C2V_COLS, BLayout::RowMajor, C2V_ROWS, C2V_COLS>;
	107	+ using GlobalFOut = GlobalTensor<
	108	+ T, pto::Shape<1, 1, 1, C2V_ROWS, C2V_COLS>, pto::Stride<TOTAL_M * N, TOTAL_M * N, C2V_ROWS * C2V_COLS, N, 1>>;
	109	+
	110	+ VecTileN vecTileHalf, tileF, tileG;
	111	+ TASSIGN(tileF, 0x10000);
	112	+ TASSIGN(tileG, 0x18000);
	113	+
	114	+ size_t fOutOffset;
	115	+ if constexpr (SplitAxis == TileSplitAxis::TILE_UP_DOWN) {
	116	+ fOutOffset = static_cast<size_t>(subBlockIdx * C2V_ROWS) * N;
	117	+ } else {
	118	+ fOutOffset = static_cast<size_t>(subBlockIdx) * C2V_COLS;
	119	+ }
	120	+ GlobalFOut globalF(srcF + fOutOffset);
	121	+ GlobalFOut globalOut(out + fOutOffset);
	122	+
	123	+ wait_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	124	+
	125	+ TPOP<PipeT, VecTileN, SplitAxis>(pipe, vecTileHalf);
	126	+ TLOAD(tileF, globalF);
	127	+
	128	+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
	129	+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
	130	+
	131	+ wait_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	132	+
	133	+ TSUB(tileG, vecTileHalf, tileF);
	134	+
	135	+ set_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	136	+
	137	+ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
	138	+ wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
	139	+
	140	+ TSTORE(globalOut, tileG);
	141	+
	142	+ set_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	143	+
	144	+ wait_flag(PIPE_V, PIPE_MTE2, EVENT_ID1);
	145	+ wait_flag(PIPE_MTE3, PIPE_V, EVENT_ID1);
	146	+
	147	+ pipe_barrier(PIPE_ALL);
	148	+}
	149	+
	150	+// ---------------------------------------------------------------------------------------
	151	+// Cube side: move both operands into L0A/L0B and matmul. Called twice with different left
	152	+// operands (from GM, then from the V2C pop), which is why it is factored out.
	153	+//
	154	+// FLAG CONTRACT: the caller must have (PIPE_FIX, PIPE_M, EVENT_ID1) and
	155	+// (PIPE_M, PIPE_MTE1, EVENT_ID1) armed. This consumes both and re-arms
	156	+// (PIPE_M, PIPE_MTE1, EVENT_ID1) for the next call or the caller's trailing drain.
	157	+// Forgetting that re-arm is what makes the core spin forever on the final wait_flag.
	158	+// ---------------------------------------------------------------------------------------
	159	+template <typename LeftT, typename RightT, typename AccT, typename SrcT, typename MatDT>
	160	+AICORE inline void movAndMatmul(LeftT& leftTile, RightT& rightTile, AccT& accTile, SrcT& srcTile, MatDT& matTileD)
	161	+{
	162	+ set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
	163	+ wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
	164	+
	165	+ wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID1);
	166	+
	167	+ TMOV(leftTile, srcTile);
	168	+ TMOV(rightTile, matTileD);
	169	+
	170	+ set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
	171	+ wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
	172	+
	173	+ wait_flag(PIPE_FIX, PIPE_M, EVENT_ID1);
	174	+
	175	+ TMATMUL(accTile, leftTile, rightTile);
	176	+
	177	+ set_flag(PIPE_M, PIPE_MTE1, EVENT_ID1);
	178	+}
	179	+
	180	+// ---------------------------------------------------------------------------------------
	181	+// Cube side.
	182	+//
	183	+// ── The ONLY difference from tpushpop_dir_both ───────────────────────────────
	184	+// There, the cube pops the V2C tile and feeds it straight into the matmul whose result it
	185	+// then pushes back C2V. That data dependency means the C2V push cannot start until the V2C
	186	+// tile has been consumed, so the two directions never hold a tile at the same time and the
	187	+// shared-slot aliasing stays hidden.
	188	+//
	189	+// Here the cube's matmul takes BOTH operands from GM, so the C2V push does not depend on
	190	+// the V2C pop, and the push is issued FIRST. Cube and vector then each have one tile in
	191	+// flight in the opposite direction simultaneously. Both are tileIndex 0, so both address
	192	+// slot 0 of GM_SLOT_BUFFER -- the two payloads land on top of each other and each consumer
	193	+// can read a mix of the two.
	194	+//
	195	+// Nothing here is exotic: the two directions have fully independent flow control (C2V uses
	196	+// FlagID/FlagID+1, V2C uses FlagID+2/FlagID+3), so neither producer ever waits on the
	197	+// other, and each is independently granted SlotNum credits against a buffer that only has
	198	+// SlotNum slots in total.
	199	+//
	200	+// BOTH directions are checked, on purpose. An earlier version of this case checked only the
	201	+// vector's output and PASSED on stock pto-isa: the two payloads are different sizes --
	202	+// C2V is a FULL slot (TOTAL_M x N x 4 = 64 KiB = SLOT_SIZE), V2C is HALF a slot
	203	+// (TOTAL_M x K x 4 = 32 KiB) -- so whichever push lands second overwrites only part of the
	204	+// other tile, and the surviving portion can be enough for one consumer to verify. The
	205	+// second matmul below feeds `popTile` into (popTile x D) and stores it, golden
	206	+// (A + B) @ D, so the cube's read is observable too. On an affected build BOTH comparisons
	207	+// fail. This is the same reason the upstream tpushpop_dir_both case passes: the defect is
	208	+// invisible unless you both (a) let the directions overlap and (b) check both directions.
	209	+// ---------------------------------------------------------------------------------------
	210	+// NOTE: these type aliases must stay INSIDE the function. Lifting them to a namespace-scope
	211	+// traits struct does not compile: C0_SIZE_BYTE is a non-dependent name there and is not
	212	+// visible at that scope in the mix-kernel TU, whereas inside the function template it is.
	213	+template <typename T, int TOTAL_M, int K, int N, TileSplitAxis SplitAxis, typename PipeT>
	214	+AICORE inline void runCubeHalf(PipeT& pipe, __gm__ T* srcA, __gm__ T* srcD, __gm__ T* outCube)
	215	+{
	216	+ constexpr uint32_t blockAlign = C0_SIZE_BYTE / sizeof(T);
	217	+ constexpr uint32_t ALIGNED_M = CeilAlign<uint32_t>(TOTAL_M, 16);
	218	+ constexpr uint32_t ALIGNED_K = CeilAlign<uint32_t>(K, blockAlign);
	219	+ constexpr uint32_t ALIGNED_N = CeilAlign<uint32_t>(N, blockAlign);
	220	+
	221	+ using PopTileV2C =
	222	+ Tile<TileType::Mat, T, ALIGNED_M, ALIGNED_K, BLayout::ColMajor, TOTAL_M, K, SLayout::RowMajor, 512>;
	223	+ using TileMatD = Tile<TileType::Mat, T, ALIGNED_K, ALIGNED_N, BLayout::ColMajor, K, N, SLayout::RowMajor, 512>;
	224	+ using GlobalD = GlobalTensor<T, pto::Shape<1, 1, 1, K, N>, pto::Stride<K * N, K * N, K * N, N, 1>>;
	225	+ using GlobalA =
	226	+ GlobalTensor<T, pto::Shape<1, 1, 1, TOTAL_M, K>, pto::Stride<TOTAL_M * K, TOTAL_M * K, TOTAL_M * K, K, 1>>;
	227	+ using GlobalOutCube =
	228	+ GlobalTensor<T, pto::Shape<1, 1, 1, TOTAL_M, N>, pto::Stride<TOTAL_M * N, TOTAL_M * N, TOTAL_M * N, N, 1>>;
	229	+
	230	+ PopTileV2C popTile, matTileA; // matTileA: left operand from GM, not from the V2C pop
	231	+ TileMatD matTileD;
	232	+ TASSIGN(matTileA, 0x20000);
	233	+ TASSIGN(matTileD, 0x40000);
	234	+
	235	+ TileLeft<T, ALIGNED_M, ALIGNED_K, TOTAL_M, K> leftTile;
	236	+ TileRight<T, ALIGNED_K, ALIGNED_N, K, N> rightTile;
	237	+ TileAcc<T, TOTAL_M, N, TOTAL_M, N> accTile;
	238	+ TASSIGN(leftTile, 0x0);
	239	+ TASSIGN(rightTile, 0x0);
	240	+ TASSIGN(accTile, 0x0);
	241	+
	242	+ GlobalA globalACube(srcA);
	243	+ GlobalD globalD(srcD);
	244	+
	245	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID1);
	246	+ set_flag(PIPE_M, PIPE_MTE1, EVENT_ID1);
	247	+
	248	+ TLOAD(matTileA, globalACube);
	249	+ TLOAD(matTileD, globalD);
	250	+
	251	+ movAndMatmul(leftTile, rightTile, accTile, matTileA, matTileD);
	252	+
	253	+ set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
	254	+ wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
	255	+
	256	+ // C2V push FIRST -- independent of anything the vector has sent.
	257	+ TPUSH<PipeT, TileAcc<T, TOTAL_M, N, TOTAL_M, N>, SplitAxis>(pipe, accTile);
	258	+
	259	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID1);
	260	+
	261	+ // V2C pop SECOND -- by now both rings hold a live tile in slot 0.
	262	+ TPOP<PipeT, PopTileV2C, SplitAxis>(pipe, popTile);
	263	+
	264	+ // Make the popped V2C tile OBSERVABLE. Without this, nothing downstream ever reads the
	265	+ // cube's V2C tile, so half the evidence is thrown away. This does NOT re-serialise the
	266	+ // directions: the C2V push above already happened.
	267	+ movAndMatmul(leftTile, rightTile, accTile, popTile, matTileD);
	268	+
	269	+ set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
	270	+ wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
	271	+
	272	+ GlobalOutCube globalOutCube(outCube);
	273	+ TSTORE(globalOutCube, accTile);
	274	+
	275	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID1);
	276	+
	277	+ wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID1);
	278	+ wait_flag(PIPE_FIX, PIPE_M, EVENT_ID1);
	279	+
	280	+ pipe_barrier(PIPE_ALL);
	281	+}
	282	+
	283	+// Computation flow (DIR_BOTH pipe). Unlike tpushpop_dir_both, the two directions OVERLAP:
	284	+// the cube's C2V push is issued BEFORE its V2C pop and does not depend on it.
	285	+// TILE_UP_DOWN: each Vec core handles TOTAL_M/2 rows, full columns
	286	+// TILE_LEFT_RIGHT: each Vec core handles full rows, K/2 or N/2 columns
	287	+//
	288	+// Vec: tileC = tileA + tileB (per vector core portion)
	289	+// Vec→Cube (V2C): TPUSH tileC → combined [TOTAL_M, K] in the FIFO
	290	+// Cube: tileE = TMATMUL(A[M,K], D[K,N]) BOTH operands from GM, so this does
	291	+// not depend on anything the vector sent
	292	+// Cube→Vec (C2V): TPUSH tileE[TOTAL_M, N] issued BEFORE the pop below
	293	+// Cube: TPOP → popTile, TSTORE TMATMUL(popTile, D) to outCube
	294	+// golden (A + B) @ D -- the cube's V2C
	295	+// read, checked so it is not silent
	296	+// Vec: TPOP → tileE_part, tileG = tileE_part - tileF, TSTORE tileG
	297	+// golden A @ D - F
	298	+template <typename T, int TOTAL_M, int K, int N, TileSplitAxis SplitAxis = TileSplitAxis::TILE_UP_DOWN>
	299	+__global__ AICORE void runTPushPopDirBoth(
	300	+ __gm__ uint64_t* ffts_addr, __gm__ T* out, __gm__ T* srcA, __gm__ T* srcB, __gm__ T* srcD, __gm__ T* srcF,
	301	+ __gm__ T* fifoMem, __gm__ T* outCube)
	302	+{
	303	+ set_ffts_base_addr((uint64_t)ffts_addr);
	304	+
	305	+ constexpr uint16_t FLAG_ID = 0;
	306	+ constexpr uint8_t FIFO_DEPTH = 2;
	307	+ constexpr uint32_t SLOT_SIZE = TOTAL_M * N * sizeof(T);
	308	+
	309	+ using BothPipe = TPipe<FLAG_ID, Direction::DIR_BOTH, SLOT_SIZE, FIFO_DEPTH>;
	310	+
	311	+ constexpr uint32_t v2cL1Base = 0x0;
	312	+ constexpr uint32_t c2vUBBase = 0x0;
	313	+
	314	+ BothPipe pipe((__gm__ void*)fifoMem, c2vUBBase, v2cL1Base);
	315	+
	316	+ if constexpr (DAV_VEC) {
	317	+ uint32_t subBlockIdx = get_subblockid();
	318	+ runVecPushV2C<T, TOTAL_M, K, N, SplitAxis, BothPipe>(pipe, srcA, srcB, subBlockIdx);
	319	+ runVecPopC2V<T, TOTAL_M, K, N, SplitAxis, BothPipe>(pipe, out, srcF, subBlockIdx);
	320	+ }
	321	+
	322	+ if constexpr (DAV_CUBE) {
	323	+ runCubeHalf<T, TOTAL_M, K, N, SplitAxis, BothPipe>(pipe, srcA, srcD, outCube);
	324	+ }
	325	+}
	326	+
	327	+template <int32_t tilingKey>
	328	+void LaunchTPushPopDirBoth(
	329	+ uint8_t* ffts, uint8_t* out, uint8_t* srcA, uint8_t* srcB, uint8_t* srcD, uint8_t* srcF, uint8_t* fifoMem,
	330	+ uint8_t* outCube, void* stream)
	331	+{
	332	+ if constexpr (tilingKey == 1) {
	333	+ runTPushPopDirBoth<float, 128, 64, 128, TileSplitAxis::TILE_UP_DOWN><<<1, nullptr, stream>>>(
	334	+ reinterpret_cast<uint64_t>(ffts), reinterpret_cast<float>(out), reinterpret_cast<float*>(srcA),
	335	+ reinterpret_cast<float>(srcB), reinterpret_cast<float>(srcD), reinterpret_cast<float*>(srcF),
	336	+ reinterpret_cast<float>(fifoMem), reinterpret_cast<float>(outCube));
	337	+ } else if constexpr (tilingKey == 2) {
	338	+ runTPushPopDirBoth<float, 128, 64, 128, TileSplitAxis::TILE_LEFT_RIGHT><<<1, nullptr, stream>>>(
	339	+ reinterpret_cast<uint64_t>(ffts), reinterpret_cast<float>(out), reinterpret_cast<float*>(srcA),
	340	+ reinterpret_cast<float>(srcB), reinterpret_cast<float>(srcD), reinterpret_cast<float*>(srcF),
	341	+ reinterpret_cast<float>(fifoMem), reinterpret_cast<float>(outCube));
	342	+ }
	343	+}
	344	+
	345	+template void LaunchTPushPopDirBoth<1>(
	346	+ uint8_t* ffts, uint8_t* out, uint8_t* srcA, uint8_t* srcB, uint8_t* srcD, uint8_t* srcF, uint8_t* fifoMem,
	347	+ uint8_t* outCube, void* stream);
	348	+template void LaunchTPushPopDirBoth<2>(
	349	+ uint8_t* ffts, uint8_t* out, uint8_t* srcA, uint8_t* srcB, uint8_t* srcD, uint8_t* srcF, uint8_t* fifoMem,
	350	+ uint8_t* outCube, void* stream);