diff --git a/Makefile b/Makefile
index ff9aa8d..bc56e79 100644
--- a/Makefile
+++ b/Makefile
@@ -27,15 +27,12 @@ check:
 	python3 test/test_tpack_element_assembly_api.py
 	python3 test/test_pr_gate_cache_order.py
 	python3 test/test_version_header.py
-	python3 test/test_layout_header_link.py
 	python3 test/test_element_tile_api.py
-	python3 test/test_tile_array_participant_mask.py
 	python3 test/test_i64_element_tile_api.py
 	python3 test/test_i64_element_tile_runtime.py
 	python3 test/test_integer_element_storage.py
 	python3 test/test_dense_f16_element_tile_api.py
 	python3 test/test_hierarchical_element_storage_api.py
-	python3 test/test_runtime_adapter_contract.py
 	$(CXX) -std=c++20 -D__linx -include test/linx_host_type_shim.hpp \
 		-fsyntax-only -Iinclude test/ptoas_linx_type_compat.cpp
 	$(CXX) -std=c++20 -D__linx -include test/linx_host_type_shim.hpp \
@@ -65,9 +62,6 @@ install:
 		echo '#define PTO_TILEOP_API_HAS_ELEMENT_TILE 1'; \
 		echo '#define PTO_TILEOP_API_HAS_ELEMENT_VIEW_METADATA 1'; \
 		echo '#define PTO_TILEOP_API_HAS_TYPED_ELEMENT_VIEWS 1'; \
-		echo '#define PTO_TILEOP_API_HAS_SYNC_CORE 1'; \
-		echo '#define PTO_TILEOP_API_HAS_GROUP_RUNTIME_ABI 1'; \
-		echo '#define PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK 1'; \
 		echo '#define PTO_TILEOP_API_HAS_512_INTEGER_ELEMENT_TILE 1'; \
 		echo '#define PTO_TILEOP_API_HAS_I64_ELEMENT_TILE_LEAF32 1'; \
 		echo '#define PTO_TILEOP_API_HAS_DENSE_F16_U16_ELEMENT_TILE 1'; \
diff --git a/docs/tileop-usage/README.md b/docs/tileop-usage/README.md
index 5df668b..ebc6fe1 100644
--- a/docs/tileop-usage/README.md
+++ b/docs/tileop-usage/README.md
@@ -30,8 +30,6 @@ kernel；其中的旧版本数量、旧 engine inventory 和旧 encoding identit
 - `PTO_TILEOP_API_REVISION`：`make install` 时实际 checkout 的 Git commit；
 - `PTO_TILEOP_API_REVISION_IS_EXACT`：revision 是否来自干净 checkout；
 - `PTO_TILEOP_API_HAS_LOCAL_B_KN_FIX`：local Right B 使用逻辑 `[K,N]` 契约。
-- `PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK`：`TileArray` raw `TMOV`
-  assembly 支持显式参与 PE mask。
 
 当前 wrapper 与 PTO 0.58.7 catalog 的覆盖边界见
 [统一支持状态表](support-status.md)。该表区分“已实现”“仅部分 layout 支持”“尚无
diff --git a/docs/tileop-usage/b-subview-b-assemble.md b/docs/tileop-usage/b-subview-b-assemble.md
index b8e10b2..b4ecc54 100644
--- a/docs/tileop-usage/b-subview-b-assemble.md
+++ b/docs/tileop-usage/b-subview-b-assemble.md
@@ -41,28 +41,6 @@ source、兼容代码和编码回归测试；`*_at_reg` 接口只出现在固定
 的 TileOP 会自动发出对应 `INIT → MIDDLE* → LAST` 生命周期的
 `B.ASSEMBLE`；`TASSEMBLY` 本身不发独立指令，只返回组装好的 parent。
 
-`TileArray<Fragment, Rows, Cols, ParticipantMask>` 的第四个模板参数指定
-raw `TMOV` assembly writer 的参与 PE 集合，默认值为 `15`。合法的非零原生
-mask 是 `1, 2, 4, 8, 12, 14, 15`；位序为 `bit = 3 - PE`，所以 PE0-only
-写作 `8`（`1000`）。同一次 INIT/MIDDLE/LAST assembly 生命周期中的所有
-`B.IOT` 都携带相同 mask。
-
-显式非默认 mask 的 writer 应放在 `PTO_PE_DOMAIN(mask)` 标记的函数中。该宏
-要求编译器提供 `__LINX_PTO_PE_DOMAIN__`；旧编译器会把调用诊断为 unavailable，
-不会静默忽略域信息。当前非默认 mask 只开放给 raw `TMOV` writer，其他仍按
-默认 mask 定义的 TileArray destination overload 会在类型检查时拒绝该 slot。
-
-```cpp
-PTO_PE_DOMAIN(8)
-void assemble_pe0(ElementTile<uint32_t, 32> (&parts)[4]) {
-  TileArray<ElementTile<uint32_t, 32>, 1, 4, 8> outputs;
-  for (int i = 0; i < 4; ++i)
-    TMOV(outputs[0][i], parts[i]);
-  auto parent = TASSEMBLY<ElementTile<uint32_t, 128>>(std::move(outputs));
-  // 在本函数域内继续消费或写回 parent。
-}
-```
-
 同一规则可以逐级组成更大的闭合 parent。每一级先完成自己的
 `INIT -> MIDDLE* -> LAST`，`TASSEMBLY` 返回的完整 `ElementTile` 才能作为下一
 级 raw `TMOV` writer 的 source。例如 F32 使用 `32 -> 128 -> 2048 -> 16384`
@@ -398,31 +376,21 @@ void session(GM &gm_lo, GM &gm_hi, Src &a, Src &b) {
 ### 完整会话示例（多 fragment 拼装一个 parent）
 
 ```cpp
-using Parent = Tile<Location::Vec, float, 32, 16, BLayout::RowMajor>;
-using GM = global_tensor<float, RowMajor<32, 4>>;
+using Dst = Tile<Location::Vec, float, 4, 8, BLayout::RowMajor>;
+using GM  = global_tensor<float, RowMajor<4, 8>>;
 
-void load_four(GM &gm0, GM &gm1, GM &gm2, GM &gm3, Parent &out) {
-  Parent tile;  // 2 KiB parent，四个 512 B writer 覆盖不同范围
-  auto first = range::assemble<4, 0>(tile);
+void load_three(GM &gm0, GM &gm1, GM &gm2, Dst &tile) {
+  auto first  = range::assemble(tile);            // INIT
   TLOAD(first, gm0);
 
-  auto middle0 = range::assemble_middle<4, 4>(tile);
-  TLOAD_ASS(middle0, gm1);
-  auto middle1 = range::assemble_middle<4, 8>(tile);
-  TLOAD_ASS(middle1, gm2);
+  auto middle = range::assemble_middle(tile);     // MIDDLE
+  TLOAD_ASS(middle, gm1);                        // 非 INIT slot 用 _ASS
 
-  auto last = range::assemble_last<4, 12>(tile);
-  TLOAD_ASS(last, gm3);
-  out = tile;  // LAST 后才发布完整 parent
+  auto last   = range::assemble_last(tile);       // LAST
+  TLOAD_ASS(last, gm2);
 }
 ```
 
-LAST 之前的 parent 处于 OPEN 状态，只能由后续 assembly slot 消费。不能
-提前 `TSTORE`、普通 `TMOV` 或通过输出引用把它写回内存；函数只发出 INIT
-后返回也不合法。使用局部 parent 完成整个 session，再向外发布。
-`AssemblyLifetime.cpp` 验证完整 session 可生成对象，并要求两种提前发布路径
-被编译器拒绝。`RangeAssemble.cpp` 中的孤立 slot 仅用于语法与 IR 检查。
-
 ## `*_ASS`：TEPL 计算直接写入 assembled destination
 
 TEPL 逐元素、归约、广播、比较接口提供 `_ASS` 后缀版本：destination 是
diff --git a/docs/tileop-usage/options.md b/docs/tileop-usage/options.md
index b090b45..1374fcb 100644
--- a/docs/tileop-usage/options.md
+++ b/docs/tileop-usage/options.md
@@ -272,55 +272,9 @@ PRelu prelu;
 TMATMUL(d, a, b, fixp::f16().prelu(prelu));
 ```
 
-## 6.1 动态有效区域与物理列数
+## 6.1 Reduction result prefix views
 
-动态 shape 的 CUBE Tile 可以保留较大的静态容量，同时让一次操作使用较小的运行时
-descriptor。`Tile<..., Rows, Cols, ..., DYNAMIC, DYNAMIC>` 中的 `Rows` 和 `Cols`
-仍决定 Tile 的容量上限和 SizeCode；它们不要求每次操作都把 `Cols` 作为 `LB2`。
-
-以下接口在原有 `UserValidCol`、`UserValidRow` 参数之后接受可选的
-`UserPhysicalCol`：
-
-- 二元逐元素操作，如 `TADD`、`TSUB`、`TMUL`、`TDIV`、位运算和 min/max；
-- `TEXPANDS`；
-- 二元 `TROWEXPAND*` 操作。
-
-三个参数都为零时，接口保留原有行为。进入运行时 geometry 重载后，省略的
-valid 维度从对应 Tile 实例读取（静态 Tile 返回类型常量），省略的物理列数使用
-`Tile::Cols`。调用时建议同时给出三个值。API 在发出 inline asm
-之前检查 `0 < ValidCol <= PhysicalCol <= Tile::Cols` 和
-`0 < ValidRow <= Tile::Rows`；非法范围直接失败。对 CUBE layout，`PhysicalCol`
-是逻辑物理列数，实际存储列数继续由 layout 和 dtype 的 CELL geometry 对齐，因此
-奇数列（例如 FP32 CUBE_M32 的 5 列）本身是合法输入。
-
-下面的 Tile 仍保留 32 KiB 的 FP32 容量，但本次计算使用
-`ValidRow=2, ValidCol=6, PhysicalCol=6`：
-
-```cpp
-using Wide = Tile<Location::Vec, float, 32, 256,
-                  BLayout::CubeM32, DYNAMIC, DYNAMIC>;
-using Row = Tile<Location::Vec, float, 32, 1,
-                 BLayout::CubeM32, DYNAMIC, 1>;
-
-Wide ga(2, 6), ba(2, 6), bf(2, 6), tmp(2, 6);
-Row mean(2);
-
-TEXPANDS(ga, 0.0f, 6, 2, 6);
-TADD(ba, ba, bf, 6, 2, 6);
-TSUB(tmp, ga, bf, 6, 2, 6);
-TROWEXPANDMUL(tmp, bf, mean, 6, 2, 6);
-```
-
-这个重载不会缩小 Tile 的分配容量，也不会把有效区拆成多个幂次宽度片段。它只让
-当前 bundle 的 `LB0/LB1/LB2` 与实际 producer descriptor 一致。CUBE `TLOAD`
-和 CUBE `TCVT` 按 valid shape 派生持久 descriptor，CUBE `TSTORE` 使用绑定源的
-descriptor；后续逐元素操作必须传递与该来源一致的运行时物理列数。完整的 6 列、
-奇数 5 列和非法范围编译 fixture 见
-`test/tileop_api/src/RuntimePhysicalCols.cpp`。
-
-## 6.2 Reduction result prefix views
-
-## 6.3 TCMPS predicate GPR carriers
+## 6.2 TCMPS predicate GPR carriers
 
 `TCMPS` 的 CUBE GPR 输出重载使用 source-only `B.IOT` 和带 destination GPR
 的 `B.IOR`，不会分配 PredicateCell。非 `U8` 类型使用低半 predicate carrier；
diff --git a/include/common/layout.hpp b/include/common/layout.hpp
index af59b78..ec5390d 100644
--- a/include/common/layout.hpp
+++ b/include/common/layout.hpp
@@ -105,7 +105,7 @@ enum PadValueEnum : uint64_t {
   Null = 3
 };
 
-inline const char *layout_type_to_str(LayoutEnum type) {
+const char *layout_type_to_str(LayoutEnum type) {
   switch (type) {
   case LayoutEnum::kRowMajor:
     return "RowMajor";
diff --git a/include/common/pto_pe_runtime.hpp b/include/common/pto_pe_runtime.hpp
deleted file mode 100644
index c8ba270..0000000
--- a/include/common/pto_pe_runtime.hpp
+++ /dev/null
@@ -1,113 +0,0 @@
-#ifndef PTO_PE_RUNTIME_HPP
-#define PTO_PE_RUNTIME_HPP
-
-// Public, compiler-facing part of the core PE rendezvous contract.  The
-// launcher owns context creation and the PE stacks; this header deliberately
-// does not describe a ready/done mailbox or an exit ABI.
-
-#include <stdint.h>
-#include <stddef.h>
-
-// The launcher passes the same entry address and descriptor to all
-// participating PEs. PE-local stack/TLS setup is complete before entry. This
-// data-only ABI is shared by compiler-rt, musl and bare launchers without
-// re-entering libc on secondary PEs.
-struct pto_runtime_group_context {
-  uint32_t magic;
-  uint16_t version;
-  uint16_t participant_count;
-  // Both point to four-element arrays indexed by architectural PE id. State
-  // publication is ordered around the runtime rendezvous.
-  volatile uint32_t *state;
-  volatile int32_t *result;
-  volatile uint32_t *error;
-};
-
-enum : uint32_t {
-  PTO_RUNTIME_GROUP_CONTEXT_MAGIC = 0x50544f47u, // "PTOG"
-  PTO_RUNTIME_GROUP_CONTEXT_VERSION = 1u,
-  PTO_RUNTIME_GROUP_STATE_READY = 0u,
-  PTO_RUNTIME_GROUP_STATE_RUNNING = 1u,
-  PTO_RUNTIME_GROUP_STATE_DONE = 2u,
-  PTO_RUNTIME_GROUP_STATE_ABORT = 3u,
-  PTO_RUNTIME_GROUP_PARTICIPANTS = 4u,
-};
-
-static_assert(sizeof(pto_runtime_group_context) == 32,
-              "PTO runtime context ABI drift");
-static_assert(offsetof(pto_runtime_group_context, state) == 8,
-              "PTO runtime context pointer layout drift");
-static_assert(offsetof(pto_runtime_group_context, result) == 16,
-              "PTO runtime context pointer layout drift");
-static_assert(offsetof(pto_runtime_group_context, error) == 24,
-              "PTO runtime context pointer layout drift");
-
-extern "C" void __linx_group_entry(const pto_runtime_group_context *)
-    __attribute__((weak, noreturn));
-extern "C" int __linx_group_pe_main(
-    unsigned, const pto_runtime_group_context *) __attribute__((weak));
-extern "C" void __pto_runtime_group_complete(
-    const pto_runtime_group_context *, int) __attribute__((weak, noreturn));
-extern "C" void __linx_group_worker_start(void)
-    __attribute__((weak, noreturn));
-
-// PTO ISA SystemRegister_THREAD_ID is SSR address 0x0026.  This is a
-// read-only architectural value; it is not the retired PEID address used by
-// older experimental launchers.
-inline uint32_t get_thread_idx() { return __builtin_linx_get_thread_idx(); }
-
-inline uint32_t get_thread_id() { return get_thread_idx(); }
-
-#ifndef PTO_TILEOP_SYNC_CORE_VIOLATION
-#define PTO_TILEOP_SYNC_CORE_VIOLATION() __builtin_trap()
-#endif
-
-#if defined(__linx) && defined(__clang__)
-#if __has_builtin(__builtin_linx_sync_core)
-#define PTO_TILEOP_COMPILER_HAS_SYNC_CORE 1
-#endif
-#endif
-
-#ifndef PTO_TILEOP_COMPILER_HAS_SYNC_CORE
-#define PTO_TILEOP_COMPILER_HAS_SYNC_CORE 0
-#endif
-
-#ifndef PTO_TILEOP_SYNC_CORE_EMIT
-#if PTO_TILEOP_COMPILER_HAS_SYNC_CORE
-#define PTO_TILEOP_SYNC_CORE_EMIT() __builtin_linx_sync_core()
-#else
-#define PTO_TILEOP_SYNC_CORE_EMIT() PTO_TILEOP_SYNC_CORE_VIOLATION()
-#endif
-#endif
-
-// Define this to 1 in a platform-specific launcher build when that launcher
-// publishes the optional context-count hook.  It is intentionally opt-in:
-// PR #957 does not standardise a guest context-count ABI.
-#ifndef PTO_TILEOP_SYNC_CORE_REQUIRE_CONTEXTS
-#define PTO_TILEOP_SYNC_CORE_REQUIRE_CONTEXTS 0
-#endif
-
-#if PTO_TILEOP_SYNC_CORE_REQUIRE_CONTEXTS
-extern "C" unsigned __pto_runtime_pe_context_count(void)
-    __attribute__((weak));
-#endif
-
-/// Rendezvous all PEs in the current core at this static site.
-///
-/// The participant set and release semantics are implemented by SYNC.CORE in
-/// the platform.  Every participating PE must execute the same site in the
-/// same dynamic order.  No hosted libc entry point is re-entered here.
-__attribute__((always_inline)) inline void SYNC_CORE() {
-#if PTO_TILEOP_SYNC_CORE_REQUIRE_CONTEXTS
-  const unsigned contexts = __pto_runtime_pe_context_count
-                                ? __pto_runtime_pe_context_count()
-                                : 0;
-  if (contexts != 4 || get_thread_idx() >= 4) {
-    PTO_TILEOP_SYNC_CORE_VIOLATION();
-    return;
-  }
-#endif
-  PTO_TILEOP_SYNC_CORE_EMIT();
-}
-
-#endif
diff --git a/include/common/pto_tile_region.hpp b/include/common/pto_tile_region.hpp
index bbe6d02..79c9ff0 100644
--- a/include/common/pto_tile_region.hpp
+++ b/include/common/pto_tile_region.hpp
@@ -47,28 +47,11 @@ inline constexpr int local_layout_code_v =
 
 // TMOV permits only the Layout attribute; its pad union is reserved-zero.
 // Zero here is an encoding requirement, not a request to fill copied padding.
-#define PTO_TMOV_LAYOUT_ASM                                                   \
-  ".if %c[ElemLayout] == 29\nB.DATR CUBE_M32, Zero\n"                        \
-  ".elseif %c[ElemLayout] == 31\nB.DATR CUBE_M16, Zero\n"                    \
+#define PTO_TMOV_LAYOUT_ASM                                                  \
+  ".if %c[ElemLayout] == 29\nB.DATR CUBE_M32, Zero\n"                       \
+  ".elseif %c[ElemLayout] == 31\nB.DATR CUBE_M16, Zero\n"                   \
   ".endif\n"
 
-#define PTO_PE_MASK_ASM(PREFIX, SUFFIX)                                      \
-  ".if %c[PEMask] == 1\n" PREFIX "0001" SUFFIX                              \
-  ".elseif %c[PEMask] == 2\n" PREFIX "0010" SUFFIX                          \
-  ".elseif %c[PEMask] == 4\n" PREFIX "0100" SUFFIX                          \
-  ".elseif %c[PEMask] == 8\n" PREFIX "1000" SUFFIX                          \
-  ".elseif %c[PEMask] == 12\n" PREFIX "1100" SUFFIX                        \
-  ".elseif %c[PEMask] == 14\n" PREFIX "1110" SUFFIX                        \
-  ".elseif %c[PEMask] == 15\n" PREFIX "1111" SUFFIX                        \
-  ".endif\n"
-
-#if defined(__LINX_PTO_PE_DOMAIN__)
-#define PTO_PE_DOMAIN(Mask) \
-  __attribute__((annotate("pto.pe_domain", Mask), noinline))
-#else
-#define PTO_PE_DOMAIN(Mask) \
-  __attribute__((unavailable("PTO_PE_DOMAIN requires compiler support for __LINX_PTO_PE_DOMAIN__")))
-#endif
 namespace region {
 
 template <typename...>
@@ -394,17 +377,14 @@ class BorrowedTileArray {
 };
 
 
-template <typename SubTile, unsigned ParticipantMask = 15>
+template <typename SubTile>
 class TileArrayOutputRef;
 
-template <typename SubTile, int Rows, int Cols, unsigned ParticipantMask = 15>
+template <typename SubTile, int Rows, int Cols>
 class TileArray {
 public:
   static_assert(Rows > 0 && Cols > 0,
                 "TileArray extents must be positive");
-  static_assert(is_valid_pe_mask(ParticipantMask) && ParticipantMask != 0,
-                "TileArray ParticipantMask must be a nonzero native PE mask");
-  static constexpr unsigned participant_mask = ParticipantMask;
   using SubTileType = SubTile;
   static constexpr int rows = Rows;
   static constexpr int cols = Cols;
@@ -454,7 +434,7 @@ class TileArray {
   public:
     Row(TileArray &array, int row) : array_(&array), row_(row) {}
 
-    TileArrayOutputRef<SubTile, ParticipantMask> operator[](int col) {
+    TileArrayOutputRef<SubTile> operator[](int col) {
       return array_->makeOutputRef(row_, col);
     }
 
@@ -467,20 +447,19 @@ class TileArray {
   Row row(int row) { return Row(*this, row); }
 
 private:
-  TileArrayOutputRef<SubTile, ParticipantMask> makeOutputRef(int row, int col) {
-    return TileArrayOutputRef<SubTile, ParticipantMask>(data_, row, col, Cols,
+  TileArrayOutputRef<SubTile> makeOutputRef(int row, int col) {
+    return TileArrayOutputRef<SubTile>(data_, row, col, Cols,
                                        Rows * Cols, ParentSizeCode);
   }
   ParentCarrier data_;
 };
 
-template <typename SubTile, unsigned ParticipantMask>
+template <typename SubTile>
 class TileArrayOutputRef {
 public:
   using SubTileType = SubTile;
-  static constexpr unsigned participant_mask = ParticipantMask;
 
-  template <typename, int, int, unsigned>
+  template <typename, int, int>
   friend class TileArray;
 
   int row() const { return row_; }
@@ -537,7 +516,7 @@ class TileArrayOutputRef {
   int parent_size_code() const { return parent_size_code_; }
 
 private:
-  template <typename, int, int, unsigned>
+  template <typename, int, int>
   friend class TileArray;
   template <typename Carrier>
   TileArrayOutputRef(Carrier &array, int row, int col, int array_cols,
@@ -561,11 +540,11 @@ template <typename Parent, typename SubTile, int Rows, int Cols>
 using BorrowedTileArray =
     region::BorrowedTileArray<Parent, SubTile, Rows, Cols>;
 
-template <typename SubTile, int Rows, int Cols, unsigned ParticipantMask = 15>
-using TileArray = region::TileArray<SubTile, Rows, Cols, ParticipantMask>;
+template <typename SubTile, int Rows, int Cols>
+using TileArray = region::TileArray<SubTile, Rows, Cols>;
 
-template <typename SubTile, unsigned ParticipantMask = 15>
-using TileArrayOutputRef = region::TileArrayOutputRef<SubTile, ParticipantMask>;
+template <typename SubTile>
+using TileArrayOutputRef = region::TileArrayOutputRef<SubTile>;
 
 template <typename T>
 struct is_subtile_view : std::false_type {};
@@ -602,8 +581,8 @@ concept is_reduction_prefix_view_v = is_reduction_prefix_view<T>::value;
 
 template <typename T>
 struct is_tile_array_output_ref : std::false_type {};
-template <typename SubTile, unsigned ParticipantMask>
-struct is_tile_array_output_ref<region::TileArrayOutputRef<SubTile, ParticipantMask>>
+template <typename SubTile>
+struct is_tile_array_output_ref<region::TileArrayOutputRef<SubTile>>
     : std::true_type {};
 template <typename T>
 inline constexpr bool is_tile_array_output_ref_v =
@@ -943,9 +922,8 @@ decltype(auto) TPARTELEMENT(const TileType &tile) {
 
 
 
-template <typename Parent, typename SubTile, int Rows, int Cols,
-          unsigned ParticipantMask>
-auto TASSEMBLY(region::TileArray<SubTile, Rows, Cols, ParticipantMask> &&array) -> Parent {
+template <typename Parent, typename SubTile, int Rows, int Cols>
+auto TASSEMBLY(region::TileArray<SubTile, Rows, Cols> &&array) -> Parent {
   static_assert(std::is_same_v<typename Parent::DType,
                                typename SubTile::DType>,
                 "TASSEMBLY parent and fragment dtypes must match");
@@ -968,7 +946,7 @@ auto TASSEMBLY(region::TileArray<SubTile, Rows, Cols, ParticipantMask> &&array)
 #ifdef __linx
   static_assert(std::is_same_v<typename Parent::TileDType,
                                typename region::TileArray<SubTile, Rows,
-                                                          Cols, ParticipantMask>::ParentCarrier>,
+                                                          Cols>::ParentCarrier>,
                 "TASSEMBLY parent and TileArray carriers must match");
   result.assignData(array.carrier());
 #endif
diff --git a/include/common/pto_tile_region_inline_asm.hpp b/include/common/pto_tile_region_inline_asm.hpp
index 6dfe147..046d4ab 100644
--- a/include/common/pto_tile_region_inline_asm.hpp
+++ b/include/common/pto_tile_region_inline_asm.hpp
@@ -692,46 +692,45 @@ PTO_REGION_SCALAR_SUBVIEW_DEST_WRAPPER(TMINS, 44)
 #undef PTO_REGION_SCALAR_DEST_WRAPPER
 #undef PTO_REGION_SCALAR_SUBVIEW_DEST_WRAPPER
 
-// Raw TileArray assembly transport. INIT allocates the parent; continuation
-// writers retain that ParentRef and use the same participant mask.
-template <int ParentSize, bool Init, bool Last, typename SubTile,
-          unsigned ParticipantMask>
+// Raw assembly transport: preserve floating payloads (including NaNs) without
+// routing a copy through numeric conversion. The fragment owns the descriptor;
+// INIT alone allocates the parent and later slots retain its exact ParentRef.
+template <int ParentSize, bool Init, bool Last, typename SubTile>
 PTO_REGION_ALWAYS_INLINE void pto_region_tmov_assemble(
-    region::TileArrayOutputRef<SubTile, ParticipantMask> &dst,
-    const SubTile &src) {
+    region::TileArrayOutputRef<SubTile> &dst, const SubTile &src) {
+  static_assert(is_element_tile_v<SubTile>,
+                "TMOV assembly currently requires an ElementTile fragment");
   static_assert(is_element_assembly_fragment_v<SubTile>,
                 "TMOV assembly requires a supported full-storage element "
                 "fragment");
   const uintptr_t offset = dst.range_base_units();
 #define PTO_REGION_TMOV_PREFIX                                               \
   "BSTART.TLSU TMOV, %D[DataType]\n"                                      \
-  PTO_TMOV_LAYOUT_ASM                                                       \
-  "B.DIM zero, %c[ValidCol], ->lb0\n"                                      \
-  "B.DIM zero, %c[ValidRow], ->lb1\n"                                      \
+  PTO_TMOV_LAYOUT_ASM                                         \
+  "B.DIM zero, %c[ValidCol], ->lb0\n"                                     \
+  "B.DIM zero, %c[ValidRow], ->lb1\n"                                     \
   "B.DIM zero, %c[PhysicalCol], ->lb2\n"
-#define PTO_REGION_TMOV_ASSEMBLE                                            \
+#define PTO_REGION_TMOV_ASSEMBLE                                             \
   "B.ASSEMBLE %c[Init], %c[Last], %[Offset], 0, %c[WriterSize]\n"
-#define PTO_REGION_TMOV_INPUTS                                              \
+#define PTO_REGION_TMOV_INPUTS                                               \
   [Src] "Tr"(src.data()),                                                   \
-  [DataType] "i"(type_traits<typename SubTile::DType>::TypeCode),           \
-  [ValidCol] "i"(SubTile::ValidCol), [ValidRow] "i"(SubTile::ValidRow),    \
-  [PhysicalCol] "i"(SubTile::Cols), [ParentSize] "i"(ParentSize),          \
-  [WriterSize] "i"(SubTile::TilesizeCode), [Offset] "r"(offset),           \
-  [Init] "i"(Init), [Last] "i"(Last), [PEMask] "i"(ParticipantMask),      \
+  [DataType] "i"(type_traits<typename SubTile::DType>::TypeCode),             \
+  [ValidCol] "i"(SubTile::ValidCol), [ValidRow] "i"(SubTile::ValidRow),        \
+  [PhysicalCol] "i"(SubTile::Cols), [ParentSize] "i"(ParentSize),             \
+  [WriterSize] "i"(SubTile::TilesizeCode), [Offset] "r"(offset),              \
+  [Init] "i"(Init), [Last] "i"(Last),                                        \
   [ElemLayout] "i"(local_layout_code_v<SubTile>)
   if constexpr (Init) {
     asm volatile(PTO_REGION_TMOV_PREFIX
-                 PTO_PE_MASK_ASM(
-                     "B.IOT %[Src].reuse, mask=",
-                     ", last, ->%[Dst]<%Z[ParentSize]>\n")
+                 "B.IOT %[Src], mask=1111, last, ->%[Dst]<%Z[ParentSize]>\n"
                  PTO_REGION_TMOV_ASSEMBLE
                  : [Dst] "=Tr"(dst.template parent_data<ParentSize>())
                  : PTO_REGION_TMOV_INPUTS
                  : "memory");
   } else {
     asm volatile(PTO_REGION_TMOV_PREFIX
-                 PTO_PE_MASK_ASM("B.IOT %[Src].reuse, mask=", ", last\n")
-                 PTO_PE_MASK_ASM("B.IOT %[Dst].reuse, mask=", "\n")
+                 "B.IOT %[Src], mask=1111, last\n"
+                 "B.IOT %[Dst], mask=1111\n"
                  PTO_REGION_TMOV_ASSEMBLE
                  :
                  : [Dst] "Tr"(dst.template parent_data<ParentSize>()),
@@ -743,10 +742,9 @@ PTO_REGION_ALWAYS_INLINE void pto_region_tmov_assemble(
 #undef PTO_REGION_TMOV_PREFIX
 }
 
-template <int ParentSize, typename SubTile, unsigned ParticipantMask>
+template <int ParentSize, typename SubTile>
 PTO_REGION_ALWAYS_INLINE void pto_region_tmov_phase(
-    region::TileArrayOutputRef<SubTile, ParticipantMask> &dst,
-    const SubTile &src) {
+    region::TileArrayOutputRef<SubTile> &dst, const SubTile &src) {
   if (dst.slot_count() == 1)
     pto_region_tmov_assemble<ParentSize, true, true>(dst, src);
   else if (dst.ordinal() == 0)
@@ -757,11 +755,10 @@ PTO_REGION_ALWAYS_INLINE void pto_region_tmov_phase(
     pto_region_tmov_assemble<ParentSize, false, false>(dst, src);
 }
 
-template <typename SubTile, unsigned ParticipantMask>
+template <typename SubTile>
   requires(is_element_assembly_fragment_v<SubTile>)
-PTO_REGION_ALWAYS_INLINE void TMOV(
-    region::TileArrayOutputRef<SubTile, ParticipantMask> dst,
-    const SubTile &src) {
+PTO_REGION_ALWAYS_INLINE void TMOV(region::TileArrayOutputRef<SubTile> dst,
+                                  const SubTile &src) {
   switch (dst.parent_size_code()) {
 #define PTO_REGION_TMOV_PARENT_CASE(N)                                      \
   case N: pto_region_tmov_phase<N>(dst, src); break
@@ -952,21 +949,19 @@ template <int ParentSize, bool Init, bool Last, typename SubTile,
 PTO_REGION_ALWAYS_INLINE void pto_region_tcvt_subview_assemble(
     region::TileArrayOutputRef<SubTile> &dst,
     region::SubTileView<Parent, SourceSubTile> &src) {
-  static_assert((SourceSubTile::BFractal == BLayout::CubeM16 ||
-                 SourceSubTile::BFractal == BLayout::CubeM32) &&
-                    SubTile::BFractal == SourceSubTile::BFractal,
-                "TCVT B.SUBVIEW/B.ASSEMBLE must preserve the CUBE_M16/M32 "
-                "source layout");
+  static_assert(SubTile::BFractal == BLayout::RowMajor &&
+                    SourceSubTile::IsCubeLayout,
+                "B.ASSEMBLE destination must be RowMajor and B.SUBVIEW "
+                "source must use a CUBE layout");
   static_assert(SubTile::SFractal == SLayout::NoneBox &&
                     SourceSubTile::SFractal == SLayout::NoneBox,
                 "inline Tile region path requires unboxed fragments");
+  static_assert(SubTile::Rows == SourceSubTile::Rows &&
+                    SubTile::Cols == SourceSubTile::Cols,
+                "TCVT TileArray slot requires matching physical shape");
   static_assert(SubTile::ValidRow == SourceSubTile::ValidRow &&
                     SubTile::ValidCol == SourceSubTile::ValidCol,
                 "TCVT TileArray slot requires matching valid shape");
-  static_assert(is_legal_tcvt_datatype_pair(
-                    type_traits<typename SourceSubTile::DType>::TypeCode,
-                    type_traits<typename SubTile::DType>::TypeCode),
-                "Illegal TCVT datatype pair");
   // PTO-ISA #265 (issue #702): field 5 is the writer extent in every phase.
   constexpr int writer_size =
       tile_type_traits<typename SubTile::TileDType>::TilesizeCode;
@@ -981,24 +976,27 @@ PTO_REGION_ALWAYS_INLINE void pto_region_tcvt_subview_assemble(
 #define PTO_REGION_TCVT_SUBVIEW_ASSEMBLY_BODY                               \
   "BSTART.TEPL 27, %D1\n"                                                  \
   "B.DATR %D2, RNONE\n"                                                    \
-  "B.DIM zero, %c4, ->lb0\n"                                               \
-  "B.DIM zero, %c5, ->lb1\n"                                               \
-  "B.IOT %3, mask=1111, last, ->%0<%Z6>\n"                                 \
-  "B.SUBVIEW 0, %7, 0, %c8\n"                                             \
-  "B.ASSEMBLE %c11, %c12, %9, 0, %c10\n"
+  "B.DIM zero, %c4, ->lb0\n"                                                   \
+  "B.DIM zero, %c5, ->lb1\n"                                                   \
+  "B.DIM zero, %c6, ->lb2\n"                                               \
+  "B.IOT %3, mask=1111, last, ->%0<%Z7>\n"                                 \
+  "B.SUBVIEW 0, %8, 0, %c9\n"                                             \
+  "B.ASSEMBLE %c12, %c13, %10, 0, %c11\n"
 #define PTO_REGION_TCVT_SUBVIEW_CONT_BODY                                   \
   "BSTART.TEPL 27, %D1\n"                                                  \
   "B.DATR %D2, RNONE\n"                                                    \
   "B.DIM zero, %c4, ->lb0\n"                                                   \
   "B.DIM zero, %c5, ->lb1\n"                                                   \
+  "B.DIM zero, %c6, ->lb2\n"                                               \
   "B.IOT %3, %0, mask=1111, last\n"                                       \
-  "B.SUBVIEW 0, %7, 0, %c8\n"                                             \
-  "B.ASSEMBLE %c11, %c12, %9, 0, %c10\n"
+  "B.SUBVIEW 0, %8, 0, %c9\n"                                             \
+  "B.ASSEMBLE %c12, %c13, %10, 0, %c11\n"
 #define PTO_REGION_TCVT_SUBVIEW_ASSEMBLY_INPUTS                            \
   "i"(type_traits<typename SourceSubTile::DType>::TypeCode),               \
   "i"(type_traits<typename SubTile::DType>::TypeCode), "Tr"(src.data()),  \
   "i"(std::remove_reference_t<decltype(src)>::ValidCol),                    \
   "i"(std::remove_reference_t<decltype(src)>::ValidRow),                    \
+  "i"(SourceSubTile::Cols),                                                 \
   "i"(ParentSize),                                                           \
   "r"(source_base_units),                                                   \
   "i"(tile_type_traits<typename SourceSubTile::TileDType>::TilesizeCode),   \
diff --git a/include/common/pto_tileop.hpp b/include/common/pto_tileop.hpp
index 1df4c85..1c221a9 100644
--- a/include/common/pto_tileop.hpp
+++ b/include/common/pto_tileop.hpp
@@ -2,7 +2,6 @@
 #define PTO_TILEOP_HPP
 
 #include "common/pto_tileop_api_revision.hpp"
-#include "common/pto_pe_runtime.hpp"
 #include "common/pto_tile.hpp"
 #include "common/pto_tile_region.hpp"
 #include "common/tileop_api.hpp"
@@ -14,4 +13,16 @@
 #include "common/pto_tile_region_inline_asm.hpp"
 #endif
 
+// Returns the current PE ID (0..3).
+// Aligned with website manual get_thread_idx(). Lowered to SSR_GET reading
+// the read-only PEID SSR (0x0802).
+inline uint32_t get_thread_idx() {
+  return __builtin_linx_get_thread_idx();
+}
+
+// API convenience alias; this is not a separate architectural operation.
+inline uint32_t get_thread_id() {
+  return get_thread_idx();
+}
+
 #endif
diff --git a/include/common/pto_tileop_api_revision.hpp b/include/common/pto_tileop_api_revision.hpp
index dc9a164..0cf926a 100644
--- a/include/common/pto_tileop_api_revision.hpp
+++ b/include/common/pto_tileop_api_revision.hpp
@@ -15,9 +15,6 @@
 #define PTO_TILEOP_API_HAS_ELEMENT_TILE 1
 #define PTO_TILEOP_API_HAS_ELEMENT_VIEW_METADATA 1
 #define PTO_TILEOP_API_HAS_TYPED_ELEMENT_VIEWS 1
-#define PTO_TILEOP_API_HAS_SYNC_CORE 1
-#define PTO_TILEOP_API_HAS_GROUP_RUNTIME_ABI 1
-#define PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK 1
 #define PTO_TILEOP_API_HAS_512_INTEGER_ELEMENT_TILE 1
 #define PTO_TILEOP_API_HAS_I64_ELEMENT_TILE_LEAF32 1
 #define PTO_TILEOP_API_HAS_INTEGER_ELEMENT_STORAGE_1024 1
diff --git a/include/jcore/template_asm.hpp b/include/jcore/template_asm.hpp
index a033768..0b3727d 100644
--- a/include/jcore/template_asm.hpp
+++ b/include/jcore/template_asm.hpp
@@ -13,6 +13,27 @@ using namespace pto;
 // the caller (they must never cross the ordinary C++ ABI as integer/stack).
 #define PTO_SHARED_INLINE __attribute__((always_inline)) inline
 
+// B.IOT/B.IOS mask syntax is a four-digit binary token, while public APIs use
+// the compact numeric PEMask template argument.  Select one canonical token
+// at assembly time without emitting a decimal mask spelling.
+#define PTO_PE_MASK_ASM(PREFIX, SUFFIX)                                        \
+  ".if %c[PEMask] == 1\n" PREFIX "0001" SUFFIX                             \
+  ".elseif %c[PEMask] == 2\n" PREFIX "0010" SUFFIX                         \
+  ".elseif %c[PEMask] == 3\n" PREFIX "0011" SUFFIX                         \
+  ".elseif %c[PEMask] == 4\n" PREFIX "0100" SUFFIX                         \
+  ".elseif %c[PEMask] == 5\n" PREFIX "0101" SUFFIX                         \
+  ".elseif %c[PEMask] == 6\n" PREFIX "0110" SUFFIX                         \
+  ".elseif %c[PEMask] == 7\n" PREFIX "0111" SUFFIX                         \
+  ".elseif %c[PEMask] == 8\n" PREFIX "1000" SUFFIX                         \
+  ".elseif %c[PEMask] == 9\n" PREFIX "1001" SUFFIX                         \
+  ".elseif %c[PEMask] == 10\n" PREFIX "1010" SUFFIX                       \
+  ".elseif %c[PEMask] == 11\n" PREFIX "1011" SUFFIX                       \
+  ".elseif %c[PEMask] == 12\n" PREFIX "1100" SUFFIX                        \
+  ".elseif %c[PEMask] == 13\n" PREFIX "1101" SUFFIX                       \
+  ".elseif %c[PEMask] == 14\n" PREFIX "1110" SUFFIX                        \
+  ".elseif %c[PEMask] == 15\n" PREFIX "1111" SUFFIX                        \
+  ".endif\n"
+
 #define PTO_RMODE_DATR_ASM(SUFFIX)                                             \
   ".if %c[RMode] == 2\nB.DATR %D[__pto_DstType], RTZ" SUFFIX                    \
   ".elseif %c[RMode] == 3\nB.DATR %D[__pto_DstType], RTM" SUFFIX                \
@@ -4168,12 +4189,12 @@ asm volatile(
 // participation 1111, no B.IOT/B.IOS members. Omitted LB0/LB1 default to one
 // and omitted LB2 to the resolved ValidCol; we pass the caller's valid shape
 // and the GM row length (logical elements) through B.DIM, and the GM base +
-// logical element row stride through B.IOR. PTO TPREFETCH does not use bytes.
+// the byte row stride (SuperScalarModel 2d467114 contract) through B.IOR.
 template <is_global_data_v gm_shape>
 void TPREFETCH(const gm_shape &src, uint32_t valid_col, uint32_t valid_row) {
-  const size_t rowStride = static_cast<size_t>(src.GetStride(3));
+  const size_t rowStride = src.GetStrideBytes(3);
   const size_t physicalCol =
-      gm_shape::Cols == DYNAMIC ? rowStride : static_cast<size_t>(gm_shape::Cols);
+      gm_shape::Cols == DYNAMIC ? rowStride : gm_shape::Cols;
   // LB2 carries the GM row length. When the GM shape is statically known it
   // lowers to an immediate (C.B.DIMI); only dynamic GM shapes keep the
   // register form. LB0/LB1 are runtime function arguments by contract.
@@ -11101,48 +11122,40 @@ PTO_SHARED_INLINE void emit_binary_assemble_init(D &dst, A &src0, B &src1) {
       : "memory");
 }
 
-// Optional runtime dimensions let a caller narrow the descriptor carried by
-// a capacity-sized Tile without changing that Tile's SizeCode. Zero keeps the
-// corresponding valid dimension from the Tile instance (or static Cols for
-// PhysicalCol); a nonzero PhysicalCol is bounded by Cols
-// and must contain ValidCol. CUBE storage rounds that logical physical column
-// through the ISA CELL geometry, so odd values remain valid where the dtype
-// and layout support them.
+// The optional ValidCol/ValidRow arguments let a caller narrow the B.DIM
+// geometry of the emitted bundle without changing the tile types: 0 (the
+// default) emits the type-derived dimension; N > 0 overrides it via the
+// runtime register form. Constraints: both values nonzero and within the
+// tile capacity when overridden; a nonzero ValidRow requires a nonzero
+// ValidCol because LB0 is mandatory in every elementwise bundle.
 
 template <is_tile_data_v T>
-constexpr bool user_dims_legal(unsigned Col, unsigned Row,
-                               unsigned PhysicalCol) {
-  if (Col == 0 || Row == 0 || PhysicalCol == 0)
-    return false;
-  if (Col > PhysicalCol || PhysicalCol > (unsigned)T::Cols ||
-      Row > (unsigned)T::Rows)
+constexpr bool user_dims_legal(unsigned Col, unsigned Row) {
+  if (Col == 0 || Row == 0)
     return false;
-  return true;
+  return Col <= (unsigned)T::Cols && Row <= (unsigned)T::Rows;
 }
 
 // One elementwise bundle: header + layout + 3x B.DIM + record + operands.
-// Runtime dimensions use register emission for lb0/lb1/lb2.
+// DimSrc selects per-slot immediate (type-derived) vs register (user)
+// emission for lb0/lb1; lb2 (physical Cols) always uses the type value.
 template <is_tile_data_v Tile, int Opcode, is_tile_data_v Src0 = Tile,
           is_tile_data_v Src1 = Tile>
 PTO_SHARED_INLINE void emit_binary_user(
-    Tile &dst, Src0 &src0, Src1 &src1, unsigned UserCol, unsigned UserRow,
-    unsigned UserPhysicalCol = 0) {
-  unsigned Col = UserCol ? UserCol : (unsigned)dst.GetValidCol();
-  unsigned Row = UserRow ? UserRow : (unsigned)dst.GetValidRow();
-  unsigned PhysicalCol = UserPhysicalCol ? UserPhysicalCol
-                                         : (unsigned)Tile::Cols;
-  if (!user_dims_legal<Tile>(Col, Row, PhysicalCol))
-    __builtin_trap();
+    Tile &dst, Src0 &src0, Src1 &src1, unsigned UserCol, unsigned UserRow) {
+  unsigned Col = UserCol ? UserCol : (unsigned)Tile::ValidCol;
+  unsigned Row = UserRow ? UserRow : (unsigned)Tile::ValidRow;
   asm volatile(
       "BSTART.TEPL %c[Opcode], %D[Type]\n"
       PTO_ELEMENTWISE_LAYOUT_ASM
       "B.DIM %[ucol], 0, ->lb0\n"
       "B.DIM %[urow], 0, ->lb1\n"
-      "B.DIM %[pcol], 0, ->lb2\n"
-      "B.IOT %[Src0].reuse, %[Src1].reuse, mask=1111, last, ->%[Dst]<%Z[TileSize]>\n"
+      "B.DIM zero, %c[Cols], ->lb2\n"
+      "B.IOT %[Src0], %[Src1], mask=1111, last, ->%[Dst]<%Z[TileSize]>\n"
       : [Dst] "=Tr"(dst.data())
-      : [ucol] "r"(Col), [urow] "r"(Row), [pcol] "r"(PhysicalCol),
+      : [ucol] "r"(Col), [urow] "r"(Row),
         [Type] "i"(type_traits<typename Tile::DType>::TypeCode),
+        [Cols] "i"(Tile::Cols),
         [Src0] "Tr"(src0.data()), [Src1] "Tr"(src1.data()),
         [Opcode] "i"(Opcode),
         [TileSize] "i"(tile_type_traits<typename Tile::TileDType>::TilesizeCode),
@@ -11150,59 +11163,21 @@ PTO_SHARED_INLINE void emit_binary_user(
       : "memory");
 }
 
-template <is_tile_data_v Tile>
-PTO_SHARED_INLINE void emit_fill_user(
-    Tile &dst, typename Tile::DType scalar, unsigned UserCol,
-    unsigned UserRow, unsigned UserPhysicalCol) {
-  const unsigned Col = UserCol ? UserCol : (unsigned)dst.GetValidCol();
-  const unsigned Row = UserRow ? UserRow : (unsigned)dst.GetValidRow();
-  const unsigned PhysicalCol = UserPhysicalCol ? UserPhysicalCol
-                                                : (unsigned)Tile::Cols;
-  if (!user_dims_legal<Tile>(Col, Row, PhysicalCol))
-    __builtin_trap();
-  typename Tile::DType scalarValue = scalar;
-  asm("" : "+r"(scalarValue));
-  asm volatile(
-      "BSTART.TEPL 59, %D[Type]\n"
-      PTO_ELEMENTWISE_LAYOUT_ASM
-      "B.DIM %[vcol], 0, ->lb0\n"
-      "B.DIM %[vrow], 0, ->lb1\n"
-      "B.DIM %[pcol], 0, ->lb2\n"
-      "B.IOT mask=1111, last, ->%[Dst]<%Z[TileSize]>\n"
-      "B.IOR [%[Scalar]],[]\n"
-      : [Dst] "=Tr"(dst.data())
-      : [vcol] "r"(Col), [vrow] "r"(Row), [pcol] "r"(PhysicalCol),
-        [Type] "i"(type_traits<typename Tile::DType>::TypeCode),
-        [TileSize] "i"(tile_type_traits<typename Tile::TileDType>::TilesizeCode),
-        [Scalar] "r"(scalarValue),
-        [ElemLayout] "i"(local_layout_code_v<Tile>)
-      : "memory");
-}
-
 } // namespace pto_elementwise_user_dims
 
 template <is_tile_data_v tile_shape>
   requires(!is_assemble_v<tile_shape>)
 void TADD(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if constexpr (is_assemble_v<tile_shape>) {
-    static_assert(tile_shape::INIT,
-                  "plain TADD requires an INIT assemble carrier; use TADD_ASS for MIDDLE/LAST");
-    // UserValidCol/UserValidRow are function parameters, not constant
-    // expressions, so they cannot participate in static_assert.  An
-    // assembled INIT destination uses the carrier geometry for its writer;
-    // reject unsupported runtime overrides before emitting the bundle.
-    if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0)
-      __builtin_trap();
-    pto_elementwise_user_dims::emit_binary_assemble_init(dst, src0, src1);
-    return;
-  }
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape);
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_elementwise_user_dims::emit_binary_user<tile_shape, 0>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
+  if constexpr (tile_shape::ValidCol > 0 || tile_shape::ValidRow > 0) {
+    if (UserValidCol != 0 || UserValidRow != 0) {
+      // User override: both dims take the runtime register form, matching
+      // the ISA B.DIM RegSrc contract for non-schema dims.
+      pto_elementwise_user_dims::emit_binary_user<tile_shape, 0>(
+          dst, src0, src1, UserValidCol, UserValidRow);
+      return;
+    }
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
     asm volatile(
@@ -11304,8 +11279,8 @@ void TADD(range::Assemble<Parent, ParentSizeCode, INIT, LAST, OffsetUnits,
   PTO_SHARED_INLINE void NAME##_U(tile_shape &dst, tile_shape &src0,           \
                                   tile_shape &src1) {                          \
     static_assert(                                                             \
-        pto_elementwise_user_dims::user_dims_legal<tile_shape>(               \
-            ValidCol, ValidRow, tile_shape::Cols),                            \
+        pto_elementwise_user_dims::user_dims_legal<tile_shape>(ValidCol,       \
+                                                              ValidRow),       \
         #NAME "_U user dims must be nonzero and within tile capacity");        \
     pto_elementwise_user_dims::emit_binary_user<tile_shape, OPCODE>(                   dst, src0, src1, ValidCol, ValidRow);                                  \
   }
@@ -11325,11 +11300,10 @@ PTO_DEFINE_ELEMENTWISE_USER_DIMS(TMIN, 12)
 // TSUB: dst = src0 - src1
 template <is_tile_data_v tile_shape>
 void TSUB(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 1>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11415,8 +11389,7 @@ void TSUB(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // backing so no payload copy or descriptor retag is introduced.
 template <is_tile_data_v Dst, is_tile_data_v Src0, is_tile_data_v Src1>
 void TMUL(Dst &dst, Src0 &src0, Src1 &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
   static_assert(std::is_same_v<typename Dst::DType, typename Src0::DType> &&
                     std::is_same_v<typename Dst::DType, typename Src1::DType>,
                 "TMUL operands must share the operation dtype");
@@ -11428,9 +11401,9 @@ void TMUL(Dst &dst, Src0 &src0, Src1 &src1,
                     Dst::ValidCol == Src1::ValidCol,
                 "TMUL operands must share layout and valid shape");
   using tile_shape = Dst;
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 2, Src0, Src1>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11515,11 +11488,10 @@ void TMUL(Dst &dst, Src0 &src0, Src1 &src1,
 // TDIV: dst = src0 / src1
 template <is_tile_data_v tile_shape>
 void TDIV(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 3>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11686,11 +11658,10 @@ void TFMOD(tile_shape &dst, tile_shape &src0, tile_shape &src1) {
 // TAND: dst = src0 & src1
 template <is_tile_data_v tile_shape>
 void TAND(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 6>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11771,11 +11742,10 @@ void TAND(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TOR: dst = src0 | src1
 template <is_tile_data_v tile_shape>
 void TOR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 7>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11856,11 +11826,10 @@ void TOR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TXOR: dst = src0 ^ src1
 template <is_tile_data_v tile_shape>
 void TXOR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 8>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -11941,11 +11910,10 @@ void TXOR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TSHL: dst = src0 << src1
 template <is_tile_data_v tile_shape>
 void TSHL(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 9>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -12026,11 +11994,10 @@ void TSHL(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TSHR: dst = src0 >> src1
 template <is_tile_data_v tile_shape>
 void TSHR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 10>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -12111,11 +12078,10 @@ void TSHR(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TMAX: dst = max(src0, src1)
 template <is_tile_data_v tile_shape>
 void TMAX(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 11>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -12196,11 +12162,10 @@ void TMAX(tile_shape &dst, tile_shape &src0, tile_shape &src1,
 // TMIN: dst = min(src0, src1)
 template <is_tile_data_v tile_shape>
 void TMIN(tile_shape &dst, tile_shape &src0, tile_shape &src1,
-          unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-          unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
+          unsigned UserValidCol = 0, unsigned UserValidRow = 0) {
+  if (UserValidCol != 0 || UserValidRow != 0) {
     pto_elementwise_user_dims::emit_binary_user<tile_shape, 12>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
+        dst, src0, src1, UserValidCol, UserValidRow);
     return;
   }
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
@@ -15589,14 +15554,7 @@ asm volatile(
 
 // TEXPANDS: broadcast scalar into dst tile
 template <is_tile_data_v tile_shape>
-void TEXPANDS(tile_shape &dst, typename tile_shape::DType s,
-              unsigned UserValidCol = 0, unsigned UserValidRow = 0,
-              unsigned UserPhysicalCol = 0) {
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_elementwise_user_dims::emit_fill_user(
-        dst, s, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
+void TEXPANDS(tile_shape &dst, typename tile_shape::DType s) {
   if constexpr (tile_shape::ValidCol > 0 && tile_shape::ValidRow > 0) {
   // Anti-fold: keep a compile-time-constant scalar (e.g. 0) off the zero
   // register so B.IOR [zero],[] still matches an instruction.
@@ -17206,8 +17164,6 @@ void TROWEXPAND(tile_shape_out &dst, tile_shape_in &src) {
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in>
 void TROWARGMAX(tile_shape_out &dst, tile_shape_in &src) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in);
-  static_assert(std::is_same_v<typename tile_shape_out::DType, uint32_t>,
-                "TROWARGMAX destination must use U32 indices");
   // ASL (row reduction): B.DIM describes the SOURCE geometry
   // (ValidCol/ValidRow/Col); the destination is rule-derived: one
   // column, ValidRow = source.ValidRow.
@@ -17291,8 +17247,6 @@ void TROWARGMAX(tile_shape_out &dst, tile_shape_in &src) {
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in>
 void TROWARGMIN(tile_shape_out &dst, tile_shape_in &src) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in);
-  static_assert(std::is_same_v<typename tile_shape_out::DType, uint32_t>,
-                "TROWARGMIN destination must use U32 indices");
   // ASL (row reduction): B.DIM describes the SOURCE geometry
   // (ValidCol/ValidRow/Col); the destination is rule-derived: one
   // column, ValidRow = source.ValidRow.
@@ -17986,53 +17940,13 @@ void TCOLARGMIN(tile_shape_out &dst, tile_shape_in &src) {
   );  }
 }
 
-namespace pto_row_expand_user_dims {
-
-template <int Opcode, int BroadcastByteOffset, is_tile_data_v Out,
-          is_tile_data_v Src0, is_tile_data_v Src1>
-PTO_SHARED_INLINE void emit_binary(
-    Out &dst, Src0 &src0, Src1 &src1, unsigned UserValidCol,
-    unsigned UserValidRow, unsigned UserPhysicalCol) {
-  const unsigned validCol = UserValidCol ? UserValidCol
-                                         : (unsigned)src0.GetValidCol();
-  const unsigned validRow = UserValidRow ? UserValidRow
-                                         : (unsigned)src0.GetValidRow();
-  const unsigned physicalCol = UserPhysicalCol ? UserPhysicalCol
-                                                : (unsigned)Out::Cols;
-  if (!pto_elementwise_user_dims::user_dims_legal<Out>(
-          validCol, validRow, physicalCol))
-    __builtin_trap();
-  asm volatile(
-      "BSTART.TEPL %c[Opcode], %D[Type]\n"
-      PTO_ROW_EXPAND_LAYOUT_ASM
-      "B.DIM %[ValidCol], 0, ->lb0\n"
-      "B.DIM %[ValidRow], 0, ->lb1\n"
-      "B.DIM %[PhysicalCol], 0, ->lb2\n"
-      "B.IOT %[Src0].reuse, %[Src1].reuse, mask=1111, last, "
-      "->%[Dst]<%Z[TileSize]>\n"
-      : [Dst] "=Tr"(dst.data())
-      : [Type] "i"(type_traits<typename Out::DType>::TypeCode),
-        [ValidCol] "r"(validCol), [ValidRow] "r"(validRow),
-        [PhysicalCol] "r"(physicalCol), [Src0] "Tr"(src0.data()),
-        [Src1] "Tr"(src1.data()), [Opcode] "i"(Opcode),
-        [TileSize] "i"(tile_type_traits<typename Out::TileDType>::TilesizeCode),
-        [BroadcastByteOffset] "i"(BroadcastByteOffset),
-        [ElemLayout] "i"(local_layout_code_v<Out>)
-      : "memory");
-}
-
-} // namespace pto_row_expand_user_dims
-
 // TROWEXPANDADD: row broadcast add
 // src1 is a per-row scalar/byte-strip operand whose shape may differ from src0
 // (see pto/TROWEXPANDMUL.md Mode 1/2); only dtype is required to match.
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDADD(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDADD(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18045,11 +17959,6 @@ void TROWEXPANDADD(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDADD source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<69, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18167,10 +18076,7 @@ void TROWEXPANDADD(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDSUB(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDSUB(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18183,11 +18089,6 @@ void TROWEXPANDSUB(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDSUB source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<70, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18305,10 +18206,7 @@ void TROWEXPANDSUB(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDMUL(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDMUL(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18321,11 +18219,6 @@ void TROWEXPANDMUL(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDMUL source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<71, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18447,10 +18340,7 @@ void TROWEXPANDMUL(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDDIV(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDDIV(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18463,11 +18353,6 @@ void TROWEXPANDDIV(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDDIV source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<72, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18585,10 +18470,7 @@ void TROWEXPANDDIV(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDMAX(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDMAX(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18601,11 +18483,6 @@ void TROWEXPANDMAX(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDMAX source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<73, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18723,10 +18600,7 @@ void TROWEXPANDMAX(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDMIN(tile_shape_out &dst, tile_shape_in0 &src0,
-                   tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                   unsigned UserValidRow = 0,
-                   unsigned UserPhysicalCol = 0) {
+void TROWEXPANDMIN(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18739,11 +18613,6 @@ void TROWEXPANDMIN(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDMIN source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<74, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -18861,10 +18730,7 @@ void TROWEXPANDMIN(tile_shape_out &dst, tile_shape_in0 &src0,
 template <is_tile_data_v tile_shape_out, is_tile_data_v tile_shape_in0,
           is_tile_data_v tile_shape_in1,
           int BroadcastByteOffset = 0>
-void TROWEXPANDEXPDIF(tile_shape_out &dst, tile_shape_in0 &src0,
-                      tile_shape_in1 &src1, unsigned UserValidCol = 0,
-                      unsigned UserValidRow = 0,
-                      unsigned UserPhysicalCol = 0) {
+void TROWEXPANDEXPDIF(tile_shape_out &dst, tile_shape_in0 &src0, tile_shape_in1 &src1) {
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in0);
   PTO_NO_SUBTILE_VIEW_ASSERT(tile_shape_in1);
   PTO_ROW_EXPAND_OFFSET_CHECK(tile_shape_out, tile_shape_in1);
@@ -18877,11 +18743,6 @@ void TROWEXPANDEXPDIF(tile_shape_out &dst, tile_shape_in0 &src0,
                     tile_shape_out::BFractal == tile_shape_in1::BFractal,
                 "TROWEXPANDEXPDIF source and broadcast layouts must match the "
                 "destination layout");
-  if (UserValidCol != 0 || UserValidRow != 0 || UserPhysicalCol != 0) {
-    pto_row_expand_user_dims::emit_binary<75, BroadcastByteOffset>(
-        dst, src0, src1, UserValidCol, UserValidRow, UserPhysicalCol);
-    return;
-  }
   if constexpr (tile_shape_out::ValidCol > 0 && tile_shape_out::ValidRow > 0) {
   static_assert(tile_carrier_width_compatible(
       type_traits<typename tile_shape_in1::DType>::TypeCode,
@@ -20739,60 +20600,27 @@ PTO_SHARED_INLINE void convert(D &dst, S &src) {
                 "the _ASS emission does not attach B.SUBVIEW, so the "
                 "subview offset would be silently dropped; consume a "
                 "TPARTVIEW SubTileView through the region wrappers instead");
-  constexpr bool IsCubeMSource =
-      S::BFractal == BLayout::CubeM16 || S::BFractal == BLayout::CubeM32;
-  static_assert(IsCubeMSource
-                    ? (D::BFractal == S::BFractal &&
-                       D::ValidRow == S::ValidRow &&
-                       D::ValidCol == S::ValidCol)
-                    : (D::Rows == S::Rows && D::Cols == S::Cols),
-                "TCVT_ASS must preserve ordinary physical shape or CUBE-M "
-                "layout and valid shape");
+  static_assert(D::Rows == S::Rows && D::Cols == S::Cols,
+                "TCVT_ASS source and destination physical shapes must match");
   constexpr unsigned WriterSizeCode = D::WriterSizeCode;
-  if constexpr (IsCubeMSource) {
-    // PTO-TCVT-CONTRACT-001: CUBE_M16/M32 binds only the logical valid
-    // shape in LB0/LB1. LB2 must be absent because physical geometry is
-    // independently derived from the destination type.
-    asm volatile(
-        "BSTART.TEPL %c[Opcode], %D[SType]\n"
-        "B.DATR %D[__pto_DstType], RNONE\n"
-        "B.DIM %[Col], 0, ->lb0\n"
-        "B.DIM %[Row], 0, ->lb1\n"
-        "B.IOT %[Src], %[Dst], mask=1111, last\n"
-        "B.ASSEMBLE %c[Init], %c[Last], %[RegSrc], %c[Off], %c[WriterSize]\n"
-        :
-        : [Opcode] "i"(Opcode),
-          [SType] "i"(type_traits<typename S::DType>::TypeCode),
-          [__pto_DstType] "i"(type_traits<typename D::DType>::TypeCode),
-          [Col] "r"(src.GetValidCol()), [Row] "r"(src.GetValidRow()),
-          [Src] "Tr"(src.data()), [Dst] "Tr"(dst.data()),
-          [RegSrc] "r"(static_cast<uintptr_t>(dst.GetRangeBase())),
-          [Init] "i"(static_cast<int>(D::INIT)),
-          [Last] "i"(static_cast<int>(D::LAST)), [Off] "i"(D::OffsetUnits),
-          [WriterSize] "i"(WriterSizeCode)
-        : "memory");
-  } else {
-    asm volatile(
-        "BSTART.TEPL %c[Opcode], %D[SType]\n"
-        "B.DATR %D[__pto_DstType], RNONE\n"
-        "B.DIM %[Col], 0, ->lb0\n"
-        "B.DIM %[Row], 0, ->lb1\n"
-        "B.DIM zero, %c[Cols], ->lb2\n"
-        "B.IOT %[Src], %[Dst], mask=1111, last\n"
-        "B.ASSEMBLE %c[Init], %c[Last], %[RegSrc], %c[Off], %c[WriterSize]\n"
-        :
-        : [Opcode] "i"(Opcode),
-          [SType] "i"(type_traits<typename S::DType>::TypeCode),
-          [__pto_DstType] "i"(type_traits<typename D::DType>::TypeCode),
-          [Col] "r"(src.GetValidCol()), [Row] "r"(src.GetValidRow()),
-          [Cols] "i"(S::Cols), [Src] "Tr"(src.data()),
-          [Dst] "Tr"(dst.data()),
-          [RegSrc] "r"(static_cast<uintptr_t>(dst.GetRangeBase())),
-          [Init] "i"(static_cast<int>(D::INIT)),
-          [Last] "i"(static_cast<int>(D::LAST)), [Off] "i"(D::OffsetUnits),
-          [WriterSize] "i"(WriterSizeCode)
-        : "memory");
-  }
+  asm volatile(
+      "BSTART.TEPL %c[Opcode], %D[SType]\n"
+      "B.DATR %D[__pto_DstType], RNONE\n"
+      "B.DIM %[Col], 0, ->lb0\n"
+      "B.DIM %[Row], 0, ->lb1\n"
+      "B.DIM zero, %c[Cols], ->lb2\n"
+      "B.IOT %[Src], %[Dst], mask=1111, last\n"
+      "B.ASSEMBLE %c[Init], %c[Last], %[RegSrc], %c[Off], %c[WriterSize]\n"
+      :
+      : [Opcode] "i"(Opcode),
+        [SType] "i"(type_traits<typename S::DType>::TypeCode),
+        [__pto_DstType] "i"(type_traits<typename D::DType>::TypeCode),
+        [Col] "r"(src.GetValidCol()), [Row] "r"(src.GetValidRow()),
+        [Cols] "i"(S::Cols), [Src] "Tr"(src.data()), [Dst] "Tr"(dst.data()),
+        [RegSrc] "r"(static_cast<uintptr_t>(dst.GetRangeBase())),
+        [Init] "i"(static_cast<int>(D::INIT)), [Last] "i"(static_cast<int>(D::LAST)),
+        [Off] "i"(D::OffsetUnits), [WriterSize] "i"(WriterSizeCode)
+      : "memory");
 }
 
 template <int Opcode, is_tile_data_v D, is_tile_data_v A, is_tile_data_v B>
diff --git a/test/common/Makefile.common b/test/common/Makefile.common
index 3af82f7..125988d 100644
--- a/test/common/Makefile.common
+++ b/test/common/Makefile.common
@@ -55,16 +55,13 @@ CC_O += --sysroot=$(LINX_SYSROOT) -nostdinc++ \
         -isystem $(LINX_SYSROOT)/usr/include/c++/v1
 CC_LINK += --sysroot=$(LINX_SYSROOT)
 endif
-# Linx fixtures use the explicit SSM.PTO v1 adapter.  The compiler-rt group
-# worker executes main on all four PEs, publishes their results around two
-# SYNC.CORE rendezvous, and calls the adapter's legal SYS/ACRC0 completion
-# service.  Keep the normal C++ runtime link: only the platform CRT startup
-# files are replaced, while libc++, libc++abi, libunwind, libc, and compiler-rt
-# remain selected by clang++.
-LINX_RUNTIME_ADAPTER := $(ROOT)/test/common/gfrun_runtime_adapter.c
-COMM_SRC_FILE += $(LINX_RUNTIME_ADAPTER)
+# Linx fixtures use the repository startup so the resulting ELF has the
+# expected BSTART.STD entry sequence.  Keep the normal C++ runtime link: only
+# the platform CRT startup files are replaced, while libc++, libc++abi,
+# libunwind, libc, and compiler-rt remain selected by clang++.
+COMM_SRC_FILE += $(ROOT)/test/common/_start.s
 COMM_SRC_DIR = $(shell dirname $(COMM_SRC_FILE))
-COMM_OBJ += $(patsubst %.c, %.o, $(subst $(COMM_SRC_DIR), $(OBJ_DIR), $(COMM_SRC_FILE)))
+COMM_OBJ += $(patsubst %.s, %.o, $(subst $(COMM_SRC_DIR), $(OBJ_DIR), $(COMM_SRC_FILE)))
 CC_LINK += -nostartfiles
 endif
 
@@ -84,7 +81,6 @@ INCLUDE += -I$(ROOT)/include -I$(ROOT)/test/common -I$(ROOT)/test/kernels/src
 QEMU = /remote/lms60/c00622284/qemu/LinxBlockModel/build/qemu-linx
 
 CC_O_ALL = $(CC_O) $(CC_VER) $(CC_OPTS)
-CC_C_O_ALL = $(filter-out -std=c++%,$(CC_O_ALL))
 # $(info "TEST_ROOT is $(TEST_ROOT)")
 # $(info "ELF_HEAD is $(ELF_HEAD)")
 # $(info "OBJ_ROOT is $(OBJ_ROOT)")
@@ -114,15 +110,15 @@ $(OBJ_DIR)%.o: $(SRC_DIR)%.cpp
 
 $(OBJ_DIR)%.o: $(SRC_DIR)%.c
 	@mkdir -p $(shell dirname $@)
-	$(CC) $(CC_C_O_ALL) $(INCLUDE) $(DEFINES) $< -o $@
+	$(CC) $(CC_O_ALL) $(INCLUDE) $(DEFINES) $< -o $@
 
 $(OBJ_DIR)%.o: $(SRC_DIR)%.s
 	@mkdir -p $(shell dirname $@)
 	$(AS) $(CC_O_ALL) $(INCLUDE) $(DEFINES) $< -o $@
 
-$(OBJ_DIR)%.o: $(COMM_SRC_DIR)%.c
+$(OBJ_DIR)%.o: $(COMM_SRC_DIR)%.s
 	@mkdir -p $(shell dirname $@)
-	$(CC) $(CC_C_O_ALL) -ffreestanding -fno-builtin $(INCLUDE) $(DEFINES) $< -o $@
+	$(AS) $(CC_O_ALL) $(INCLUDE) $(DEFINES) $< -o $@
 
 $(TARGET): $(OBJ) $(COMM_OBJ)
 	@mkdir -p $(shell dirname $@)
diff --git a/test/common/_start.s b/test/common/_start.s
index 62617e5..537b98e 100644
--- a/test/common/_start.s
+++ b/test/common/_start.s
@@ -1,5 +1,3 @@
-# Legacy direct-entry fixture. It is intentionally not linked by the current
-# gfrun benchmark path; use gfrun_runtime_adapter.c for hosted completion.
 .global  _start
   .type   _start,@function
   .text
@@ -9,4 +7,4 @@ _start:
 _end:
   bstart.aux fall
   acrc 3
-  c.bstop
+  c.bstop
\ No newline at end of file
diff --git a/test/common/gfrun_runtime_adapter.c b/test/common/gfrun_runtime_adapter.c
deleted file mode 100644
index b920e63..0000000
--- a/test/common/gfrun_runtime_adapter.c
+++ /dev/null
@@ -1,51 +0,0 @@
-// Explicit SSM.PTO v1 platform adapter for repository benchmarks.
-// The compiler-rt group worker owns synchronization and result publication;
-// this file owns only model admission, the main() callback, and completion.
-
-#include <stdint.h>
-
-struct pto_runtime_group_context;
-
-extern int main(void);
-extern void __linx_group_worker_start(void) __attribute__((noreturn));
-
-// Writable slot populated by the gfrun launcher before any PE executes.
-volatile uintptr_t __ssm_pto_group_context;
-
-__attribute__((used, section(".note.ssm.pto"), aligned(4)))
-static const uint32_t ssm_hosted_note[] = {
-    8, 16, 0x50544f47, 0x2e4d5353, 0x004f5450,
-    1, 4, 32, 0,
-};
-
-const struct pto_runtime_group_context *
-__pto_runtime_group_context_get(void) {
-  return (const struct pto_runtime_group_context *)__ssm_pto_group_context;
-}
-
-int __linx_group_pe_main(unsigned pe,
-                         const struct pto_runtime_group_context *context) {
-  (void)pe;
-  (void)context;
-  return main();
-}
-
-__attribute__((noreturn)) void __pto_runtime_group_complete(
-    const struct pto_runtime_group_context *context, int status) {
-  const uintptr_t service = 0x50544f01u;
-  __asm__ volatile(
-      "BSTART.SYS FALL\n"
-      "ori %0, 0, ->a0\n"
-      "ori %1, 0, ->a1\n"
-      "ori %2, 0, ->x1\n"
-      "acrc 0\n"
-      "BSTOP\n"
-      :
-      : "r"(context), "r"((intptr_t)status), "r"(service)
-      : "a0", "a1", "x1", "memory");
-  __builtin_unreachable();
-}
-
-__attribute__((noreturn)) void _start(void) {
-  __linx_group_worker_start();
-}
diff --git a/test/test_element_tile_api.py b/test/test_element_tile_api.py
index 01dee9f..1cf5c2d 100644
--- a/test/test_element_tile_api.py
+++ b/test/test_element_tile_api.py
@@ -446,7 +446,7 @@ def test_release_rejects_bad_ranges(self):
                     run=True, release=True, runtime_reject=True)
 
     def test_reject_unsupported_dtype(self):
-        self.compile('pto::ElementTile<unsigned long,32> supported;\n')
+        self.compile('pto::ElementTile<unsigned long,32> unsupported;\n', accepts=False)
         self.compile('pto::ElementTile<int16_t,32> unsupported;\n', accepts=False)
 
     def test_reject_unsupported_capacity(self):
diff --git a/test/test_layout_header_link.py b/test/test_layout_header_link.py
deleted file mode 100644
index 951ca62..0000000
--- a/test/test_layout_header_link.py
+++ /dev/null
@@ -1,41 +0,0 @@
-#!/usr/bin/env python3
-"""A public header must link across translation units without symbol rewriting."""
-import os
-from pathlib import Path
-import shlex
-import subprocess
-import tempfile
-import unittest
-
-ROOT = Path(__file__).resolve().parents[1]
-
-
-class LayoutHeaderLinkTest(unittest.TestCase):
-    def test_inline_layout_name_has_one_address_across_translation_units(self):
-        with tempfile.TemporaryDirectory() as directory:
-            out = Path(directory)
-            for name in ("left", "right"):
-                (out / f"{name}.cpp").write_text(
-                    '#include "common/layout.hpp"\n'
-                    'using Function = const char *(*)(pto::LayoutEnum);\n'
-                    f'Function {name}() {{ return pto::layout_type_to_str; }}\n')
-            (out / "main.cpp").write_text(
-                '#include "common/layout.hpp"\n#include <cstring>\n'
-                'using Function = const char *(*)(pto::LayoutEnum);\n'
-                'Function left(); Function right();\n'
-                'int main() {\n'
-                '  if (left() != right()) return 1;\n'
-                '  if (std::strcmp(left()(pto::LayoutEnum::kRowMajor), "RowMajor")) return 2;\n'
-                '  if (std::strcmp(right()(pto::LayoutEnum::kColMajor), "ColMajor")) return 3;\n'
-                '  return 0;\n}\n')
-            command = shlex.split(os.environ.get("CXX", "c++"))
-            command += ["-std=c++20", "-O0", "-I", str(ROOT / "include")]
-            command += [str(out / f"{name}.cpp") for name in ("left", "right", "main")]
-            command += ["-o", str(out / "layout-link")]
-            result = subprocess.run(command, capture_output=True, text=True)
-            self.assertEqual(result.returncode, 0, result.stdout + result.stderr)
-            subprocess.run([str(out / "layout-link")], check=True)
-
-
-if __name__ == "__main__":
-    unittest.main()
diff --git a/test/test_runtime_adapter_contract.py b/test/test_runtime_adapter_contract.py
deleted file mode 100644
index 0c21924..0000000
--- a/test/test_runtime_adapter_contract.py
+++ /dev/null
@@ -1,53 +0,0 @@
-#!/usr/bin/env python3
-
-from pathlib import Path
-
-
-ROOT = Path(__file__).resolve().parents[1]
-ADAPTER = ROOT / "test/common/gfrun_runtime_adapter.c"
-MAKEFILE = ROOT / "test/common/Makefile.common"
-
-
-def significant_lines(text: str) -> list[str]:
-    return [
-        line.strip().lower()
-        for line in text.splitlines()
-        if line.strip() and not line.lstrip().startswith("#")
-    ]
-
-
-def main() -> None:
-    source = ADAPTER.read_text()
-    lines = significant_lines(source)
-    acrc_index = next(i for i, line in enumerate(lines) if '"acrc 0\\n"' in line)
-
-    preceding_start = next(
-        line for line in reversed(lines[:acrc_index]) if "bstart." in line
-    )
-    if "bstart.sys" not in preceding_start:
-        raise SystemExit(f"ACRC is not in a SYS block: {preceding_start}")
-    if '"bstop\\n"' not in lines[acrc_index + 1]:
-        raise SystemExit("ACRC must be the terminal scalar operation")
-    required_contract = {
-        "__linx_group_worker_start": "compiler-rt group worker entry",
-        "__linx_group_pe_main": "per-PE main callback",
-        "return main();": "benchmark result propagation",
-        ".note.ssm.pto": "explicit hosted ABI admission note",
-        "0x50544f01u": "hosted completion service tag",
-        '"a0", "a1", "x1", "memory"': "completion register contract",
-    }
-    for token, description in required_contract.items():
-        if token not in source:
-            raise SystemExit(f"runtime adapter lacks {description}: {token}")
-
-    makefile = MAKEFILE.read_text()
-    if "LINX_RUNTIME_ADAPTER := $(ROOT)/test/common/gfrun_runtime_adapter.c" not in makefile:
-        raise SystemExit("Linx fixtures do not select the runtime adapter")
-    if "COMM_SRC_FILE += $(LINX_RUNTIME_ADAPTER)" not in makefile:
-        raise SystemExit("runtime adapter is not linked by default")
-
-    print("PASS: benchmark runtime adapter uses the SYS/ACRC completion path")
-
-
-if __name__ == "__main__":
-    main()
diff --git a/test/test_sync_core_contract.cpp b/test/test_sync_core_contract.cpp
deleted file mode 100644
index 65241d4..0000000
--- a/test/test_sync_core_contract.cpp
+++ /dev/null
@@ -1,47 +0,0 @@
-#include <cstdlib>
-
-static unsigned test_context_count;
-static unsigned test_pe_id;
-static unsigned test_sync_calls;
-static unsigned test_violations;
-
-extern "C" unsigned __pto_runtime_pe_context_count(void) {
-  return test_context_count;
-}
-
-inline unsigned __builtin_linx_get_thread_idx() { return test_pe_id; }
-
-#define PTO_TILEOP_SYNC_CORE_REQUIRE_CONTEXTS 1
-#define PTO_TILEOP_SYNC_CORE_VIOLATION() (++test_violations)
-#define PTO_TILEOP_SYNC_CORE_EMIT() (++test_sync_calls)
-#include <common/pto_pe_runtime.hpp>
-
-static_assert(sizeof(pto_runtime_group_context) == 32);
-static_assert(PTO_RUNTIME_GROUP_PARTICIPANTS == 4);
-
-static void reset(unsigned contexts, unsigned pe_id) {
-  test_context_count = contexts;
-  test_pe_id = pe_id;
-  test_sync_calls = 0;
-  test_violations = 0;
-}
-
-int main() {
-  reset(1, 0);
-  SYNC_CORE();
-  if (test_sync_calls != 0 || test_violations != 1)
-    return EXIT_FAILURE;
-
-  reset(4, 4);
-  SYNC_CORE();
-  if (test_sync_calls != 0 || test_violations != 1)
-    return EXIT_FAILURE;
-
-  for (unsigned pe_id = 0; pe_id != 4; ++pe_id) {
-    reset(4, pe_id);
-    SYNC_CORE();
-    if (test_sync_calls != 1 || test_violations != 0)
-      return EXIT_FAILURE;
-  }
-  return EXIT_SUCCESS;
-}
diff --git a/test/test_tile_array_participant_mask.py b/test/test_tile_array_participant_mask.py
deleted file mode 100644
index 6cf6b55..0000000
--- a/test/test_tile_array_participant_mask.py
+++ /dev/null
@@ -1,195 +0,0 @@
-#!/usr/bin/env python3
-"""Compile-time and native contracts for TileArray participant masks."""
-
-import os
-from pathlib import Path
-import re
-import shutil
-import subprocess
-import tempfile
-import unittest
-
-
-ROOT = Path(__file__).resolve().parents[1]
-FIXTURE = ROOT / "test/tileop_api/src/TileArrayParticipantMask.cpp"
-
-
-class TileArrayParticipantMaskTest(unittest.TestCase):
-    def host_compile(self, declaration, accepts):
-        compiler = os.environ.get("CXX") or shutil.which("c++")
-        self.assertIsNotNone(compiler)
-        with tempfile.TemporaryDirectory(prefix="tile-array-mask-host-") as work:
-            source = Path(work) / "contract.cpp"
-            source.write_text(
-                "#include <common/pto_tile.hpp>\n"
-                "#include <common/pto_tile_region.hpp>\n"
-                "using namespace pto;\n" + declaration
-            )
-            result = subprocess.run([
-                compiler, "-std=c++20", "-D__linx", "-include",
-                str(ROOT / "test/linx_host_type_shim.hpp"), "-I",
-                str(ROOT / "include"), "-fsyntax-only", str(source),
-            ], text=True, capture_output=True)
-            self.assertEqual(result.returncode == 0, accepts, result.stderr)
-
-    def test_public_type_and_invalid_masks(self):
-        self.host_compile(
-            "using A=TileArray<ElementTile<unsigned,32>,1,4,8>;"
-            "static_assert(A::participant_mask==8);"
-            "using R=decltype(std::declval<A&>()[0][0]);"
-            "static_assert(R::participant_mask==8);",
-            True,
-        )
-        for mask in (1, 2, 4, 8, 12, 14, 15):
-            with self.subTest(valid_mask=mask):
-                self.host_compile(
-                    f"TileArray<ElementTile<unsigned,32>,1,4,{mask}> ok;",
-                    True,
-                )
-        for mask in (0, 3, 5, 16):
-            with self.subTest(mask=mask):
-                self.host_compile(
-                    f"TileArray<ElementTile<unsigned,32>,1,4,{mask}> bad;",
-                    False,
-                )
-
-    def test_nondefault_mask_rejects_unadapted_writer(self):
-        self.host_compile(
-            "using Leaf=ElementTile<unsigned,32>;"
-            "void default_writer(TileArrayOutputRef<Leaf>);"
-            "void bad(){"
-            "TileArray<ElementTile<unsigned,32>,1,4,8> a;"
-            "default_writer(a[0][0]);}",
-            False,
-        )
-        self.host_compile(
-            "PTO_PE_DOMAIN(8) void owned(){} int main(){ owned(); }",
-            False,
-        )
-
-    def native_common(self):
-        compiler = os.environ.get("LINX_CXX")
-        resource = os.environ.get("LINX_RESOURCE_DIR")
-        sysroot = os.environ.get("LINX_SYSROOT")
-        if not (compiler and resource and sysroot):
-            self.skipTest("set LINX_CXX/LINX_RESOURCE_DIR/LINX_SYSROOT")
-        return compiler, [
-            compiler, "-std=c++20", "--target=linx64v5-unknown-linux-musl",
-            "-D__linx", "-mlxbc", "-fenable-matrix", "-O2", "-mllvm",
-            "-enable-all-vector-as-tilereg=true", "-resource-dir=" + resource,
-            "-mllvm", "-linxv5-enable-generic-element=true",
-            "-fno-vectorize", "-fno-slp-vectorize", "-Werror=unknown-pragmas",
-            "--sysroot=" + sysroot, "-I", str(ROOT / "include"),
-        ]
-
-    def test_native_nested_pe0_and_default_masks(self):
-        compiler, common = self.native_common()
-        with tempfile.TemporaryDirectory(prefix="tile-array-mask-native-") as work:
-            work = Path(work)
-            ir = work / "mask.ll"
-            result = subprocess.run(
-                common + ["-S", "-emit-llvm", str(FIXTURE), "-o", str(ir)],
-                text=True, capture_output=True,
-            )
-            self.assertEqual(result.returncode, 0, result.stderr)
-            text = ir.read_text()
-            pe0 = text[text.index("define dso_local void @assemble_pe0"):]
-            pe0 = pe0[:pe0.index("define dso_local void @assemble_default")]
-            self.assertIn("B.ASSEMBLE", pe0)
-            self.assertIn("== 8\\0A", pe0)
-            self.assertRegex(pe0, r"i1 false, i32 8, i32 29")
-
-            obj = work / "mask.o"
-            result = subprocess.run(
-                common + ["-c", str(FIXTURE), "-o", str(obj)],
-                text=True, capture_output=True,
-            )
-            self.assertEqual(result.returncode, 0, result.stderr)
-            dumped = subprocess.run([
-                str(Path(compiler).parent / "llvm-objdump"), "-d",
-                "--disassembler-options=no-tile-macros", str(obj),
-            ], text=True, capture_output=True)
-            self.assertEqual(dumped.returncode, 0, dumped.stderr)
-            pe0_asm = dumped.stdout[
-                dumped.stdout.index("<assemble_pe0>"):
-                dumped.stdout.index("<assemble_default>")
-            ]
-            assembly_blocks = [
-                block for block in re.findall(
-                    r"BSTART(?:\.TLSU\s+|\.)TMOV\b[\s\S]*?(?=BSTART|$)", pe0_asm
-                ) if "B.ASSEMBLE" in block
-            ]
-            self.assertEqual(len(assembly_blocks), 4)
-            for block in assembly_blocks:
-                self.assertRegex(block, r"B\.IOT[^\n]*mask=1000")
-                self.assertNotRegex(block, r"B\.IOT[^\n]*mask=1111")
-            self.assertRegex(pe0_asm, r"B\.ASSEMBLE\s+1, 0, [^,]+, 0, 1")
-            self.assertRegex(pe0_asm, r"BSTART\.TEPL\s+TXOR")
-            self.assertEqual(
-                len(re.findall(r"BSTART(?:\.TLSU\s+|\.)TSTORE\b", pe0_asm)), 2
-            )
-            default_asm = dumped.stdout[dumped.stdout.index("<assemble_default"):]
-            self.assertRegex(default_asm, r"B\.IOT[^\n]*mask=1111")
-
-    def test_native_nondefault_mask_rejects_unadapted_tcvt(self):
-        _, common = self.native_common()
-        with tempfile.TemporaryDirectory(prefix="tile-array-mask-negative-") as work:
-            source = Path(work) / "bad.cpp"
-            source.write_text(
-                "#include <common/pto_tileop.hpp>\n"
-                "using namespace pto;\n"
-                "void bad(){ TileArray<ElementTile<unsigned,32>,1,4,8> a; "
-                "ElementTile<unsigned,32> x; TCVT(a[0][0], x); }\n"
-            )
-            result = subprocess.run(
-                common + ["-fsyntax-only", str(source)],
-                text=True, capture_output=True,
-            )
-            self.assertNotEqual(result.returncode, 0)
-            self.assertIn("no matching function", result.stderr)
-
-    def test_native_function_domain_annotation(self):
-        _, common = self.native_common()
-        with tempfile.TemporaryDirectory(prefix="pe-domain-annotation-") as work:
-            source = Path(work) / "domain.cpp"
-            ir = Path(work) / "domain.ll"
-            source.write_text(
-                "#include <common/pto_tile_region.hpp>\n"
-                "PTO_PE_DOMAIN(8) void owned(unsigned *p){ *p = 8; }\n"
-            )
-            result = subprocess.run(
-                common + ["-S", "-emit-llvm", str(source), "-o", str(ir)],
-                text=True, capture_output=True,
-            )
-            self.assertEqual(result.returncode, 0, result.stderr)
-            text = ir.read_text()
-            self.assertIn('c"pto.pe_domain\\00"', text)
-            self.assertRegex(text, r"@\.args = .*\{ i32 8 \}")
-
-    def test_native_rejects_tile_reference_domain_boundary(self):
-        _, common = self.native_common()
-        with tempfile.TemporaryDirectory(prefix="pe-domain-tile-boundary-") as work:
-            source = Path(work) / "bad.cpp"
-            source.write_text(
-                "#include <common/pto_tileop.hpp>\n"
-                "using namespace pto;\n"
-                "using Tile128=ElementTile<unsigned,128>;\n"
-                "__attribute__((noinline)) static void consume("
-                "Tile128 &input, unsigned *output){Tile128 result;"
-                "TADDS(result,input,1u); TSTORE(output,result,128);}\n"
-                "PTO_PE_DOMAIN(8) void bad(Tile128 &input, unsigned *output){"
-                "consume(input,output);}\n"
-            )
-            result = subprocess.run(
-                common + ["-c", str(source), "-o", str(Path(work) / "bad.o")],
-                text=True, capture_output=True,
-            )
-            self.assertNotEqual(result.returncode, 0)
-            self.assertIn(
-                "Tile-reference arguments may not cross a function PE domain boundary",
-                result.stderr,
-            )
-
-
-if __name__ == "__main__":
-    unittest.main()
diff --git a/test/test_v058_engine_contract.py b/test/test_v058_engine_contract.py
index 2701eb8..d8f69f1 100644
--- a/test/test_v058_engine_contract.py
+++ b/test/test_v058_engine_contract.py
@@ -233,40 +233,6 @@ def test_tcvt_cube_m_encodes_only_source_valid_col_and_row(self) -> None:
         self.assertIn('CubeTileM16', fixture)
         self.assertIn('CubeTileM32', fixture)
 
-    def test_tcvt_ass_omits_lb2_only_for_cube_m(self) -> None:
-        convert = re.search(
-            r'(?s)template <int Opcode, is_tile_data_v D, is_tile_data_v S>\n'
-            r'PTO_SHARED_INLINE void convert\(D &dst, S &src\) \{.*?\n\}',
-            self.header,
-        )
-        self.assertIsNotNone(convert)
-        carrier = convert.group(0)
-        self.assertIn('constexpr bool IsCubeMSource', carrier)
-        cube_branch, ordinary_branch = carrier.split(
-            'if constexpr (IsCubeMSource) {', 1)[1].split('} else {', 1)
-        self.assertNotIn('->lb2', cube_branch)
-        self.assertIn('->lb2', ordinary_branch)
-
-    def test_tcvt_subview_assembly_preserves_cube_and_omits_lb2(self) -> None:
-        region_header = (ROOT / 'include' / 'common' /
-                         'pto_tile_region_inline_asm.hpp').read_text(
-                             encoding='utf-8')
-        carrier = re.search(
-            r'(?s)void pto_region_tcvt_subview_assemble\(.*?\n\}',
-            region_header,
-        )
-        self.assertIsNotNone(carrier)
-        body = carrier.group(0)
-        self.assertIn('SubTile::BFractal == SourceSubTile::BFractal', body)
-        self.assertIn('SourceSubTile::BFractal == BLayout::CubeM16', body)
-        self.assertIn('SourceSubTile::BFractal == BLayout::CubeM32', body)
-        self.assertNotIn('->lb2', body)
-
-        fixture = (ROOT / 'test' / 'tileop_api' / 'src' /
-                   'TileRegionTCVTSubviewAssembly.cpp').read_text(
-                       encoding='utf-8')
-        self.assertIn('CubeTileM32<__bf16', fixture)
-
     def test_range_modifiers_expose_simple_factories(self) -> None:
         tile_header = PTO_TILE.read_text(encoding="utf-8")
         tile = (ROOT / "test" / "tileop_api" / "src" /
@@ -802,17 +768,20 @@ def test_static_valid_shape_bindings_remain_immediate_eligible(self) -> None:
             self.assertIn(spelling, header)
         # A register must never land in the immediate slot of B.DIM.
         self.assertNotIn("B.DIM zero, %[", header)
-        # lb2 stays immediate unless the public call explicitly supplies a
-        # runtime physical column, or TPREFETCH uses its dynamic-GM fallback.
+        # lb2 must always stay the static immediate form; the only allowed
+        # dynamic lb2 is TPREFETCH's dynamic-GM fallback branch.
         dyn_lb2 = [
             line for line in header.split("\n")
             if "->lb2" in line and "B.DIM" in line and "%c" not in line
         ]
-        self.assertEqual(len(dyn_lb2), 4)
-        self.assertTrue(any('%[pcol], 0, ->lb2' in line for line in dyn_lb2))
-        self.assertTrue(any('%[PhysicalCol], 0, ->lb2' in line
-                            for line in dyn_lb2))
-        self.assertTrue(any('%[Col], 0, ->lb2' in line for line in dyn_lb2))
+        tprefetch = header[header.index("void TPREFETCH"):]
+        tprefetch = tprefetch[: tprefetch.index("\n}")]
+        self.assertEqual(
+            len(dyn_lb2),
+            tprefetch.count('"B.DIM %[Col], 0, ->lb2\\n"'),
+            "lb2 must be immediate-form everywhere except TPREFETCH's "
+            "dynamic-GM fallback",
+        )
         # TMATMUL: M stays a runtime register (group_M), N/K are compile-time
         # immediates per resolve_matmul_shape.
         self.assertIn('"B.DIM %[M], 0, ->lb0\\n"', header)
@@ -850,16 +819,6 @@ def test_cube_accumulator_is_explicit_and_destination_is_distinct(self) -> None:
         self.assertIn('"B.IOT %[C], mask=1111\\n"', self.header)
         self.assertIn('[Dst] "=&Tr"(dst.data())', self.header)
 
-    def test_tprefetch_uses_element_stride(self) -> None:
-        tprefetch = self.header[self.header.index("void TPREFETCH"):]
-        tprefetch = tprefetch[: tprefetch.index("\n}")]
-        self.assertIn("src.GetStride(3)", tprefetch)
-        self.assertNotIn("src.GetStrideBytes(3)", tprefetch)
-
-    def test_rowarg_destinations_are_u32(self) -> None:
-        self.assertIn("TROWARGMAX destination must use U32 indices", self.header)
-        self.assertIn("TROWARGMIN destination must use U32 indices", self.header)
-
     def test_tgemv_uses_a_then_b_source_and_type_order(self) -> None:
         self.assertIn('"B.IOT %[Vec], %[Mtx], mask=1111\\n"', self.header)
         self.assertNotIn('"B.IOT %[Mtx], %[Vec], mask=1111\\n"', self.header)
diff --git a/test/test_version_header.py b/test/test_version_header.py
index d4b36e8..e61a9e7 100644
--- a/test/test_version_header.py
+++ b/test/test_version_header.py
@@ -22,9 +22,6 @@ def test_source_header_exposes_version_and_feature(self):
         self.assertIn('#define PTO_TILEOP_API_HAS_ELEMENT_TILE 1', header)
         self.assertIn('#define PTO_TILEOP_API_HAS_ELEMENT_VIEW_METADATA 1', header)
         self.assertIn('#define PTO_TILEOP_API_HAS_TYPED_ELEMENT_VIEWS 1', header)
-        self.assertIn(
-            '#define PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK 1', header
-        )
         self.assertIn(
             '#define PTO_TILEOP_API_HAS_DENSE_F16_U16_ELEMENT_TILE 1', header
         )
@@ -77,10 +74,6 @@ def test_install_records_exact_checkout_revision(self):
             self.assertIn('#define PTO_TILEOP_API_HAS_ELEMENT_TILE 1', installed)
             self.assertIn('#define PTO_TILEOP_API_HAS_ELEMENT_VIEW_METADATA 1', installed)
             self.assertIn('#define PTO_TILEOP_API_HAS_TYPED_ELEMENT_VIEWS 1', installed)
-            self.assertIn(
-                '#define PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK 1',
-                installed,
-            )
             self.assertIn(
                 '#define PTO_TILEOP_API_HAS_DENSE_F16_U16_ELEMENT_TILE 1',
                 installed,
diff --git a/test/tileop_api/compile.all b/test/tileop_api/compile.all
index 11250f8..740d912 100755
--- a/test/tileop_api/compile.all
+++ b/test/tileop_api/compile.all
@@ -78,18 +78,10 @@ run_link_smoke() {
   "$MAKE_CMD" --no-print-directory TESTCASE=TCI all
   local elf="$(cd "$(dirname "$0")/../.." && pwd)/output/tileop_api/elf/tileop_api_TCI_${PLAT}.elf"
   test -s "$elf"
-  "$COMPILER_DIR/llvm-readobj" --file-headers --sections --symbols "$elf" >"${elf}.readobj"
+  "$COMPILER_DIR/llvm-readobj" --file-headers --symbols "$elf" >"${elf}.readobj"
   grep -Eq 'Entry|Name: _start' "${elf}.readobj"
-  grep -q 'Name: __linx_group_worker_start' "${elf}.readobj"
-  grep -q 'Name: __linx_group_pe_main' "${elf}.readobj"
-  grep -q '.note.ssm.pto' "${elf}.readobj"
   "$COMPILER_DIR/llvm-objdump" -dr "$elf" >"${elf}.diss"
   grep -q 'BSTART.STD' "${elf}.diss"
-  # LLVM's canonical disassembly name for the architectural SYS/AUX block
-  # encoding is BSTART.AUX; the source contract separately requires the
-  # explicit BSTART.SYS spelling in the runtime adapter.
-  grep -Eq 'BSTART\.(SYS|AUX)' "${elf}.diss"
-  grep -q 'acrc 0' "${elf}.diss"
   echo "PASS link-smoke: $elf"
 }
 
diff --git a/test/tileop_api/run_negatives.sh b/test/tileop_api/run_negatives.sh
index 90eda1e..30f54d2 100755
--- a/test/tileop_api/run_negatives.sh
+++ b/test/tileop_api/run_negatives.sh
@@ -48,7 +48,6 @@ GMOV_CASES="fp64 s64 u64 dtype shape valid_shape layout capacity location shared
 TCVT_CASES="cube_layout cube_valid_shape cube_n8 valid_shape assemble_valid_shape same_type_rto_fp16 same_type_rto_bf16 same_type_rto_fp32 same_type_rto_fp64 same_type_rto_reinterpreted same_type_rto_prefix same_type_rto_tilearray"
 TEXPDIF_CASES="type narrow n8 mixed_layout shape"
 TCI2D_CASES="layout location dtype shape steps m16_rows"
-ROWARG_CASES="max min"
 I64_CASES="fp64_element unsupported_parent m16 short_m32 subview_pair assemble_pair writer_limit"
 TCOLSUM_CASES="dest_rows dest_cols dtype"
 ISSUE288_CASES="layout dtype shape valid_shape"
@@ -87,12 +86,6 @@ expect_rejected() {
   fi
 }
 
-for c in $ROWARG_CASES; do
-  define=SHOULD_FAIL_TROWARG$(echo "$c" | tr '[:lower:]' '[:upper:]')_NON_U32
-  expect_rejected "trowarg_$c" "$define" RowArgNegatives.cpp \
-    'TROWARG(MAX|MIN) destination must use U32 indices'
-done
-
 for c in $CASES; do
   if "$CXX" "${FLAGS[@]}" -DSHOULD_FAIL_$c src/PostProcessNegatives.cpp \
        -o "$OUT/neg_$c.o" >/dev/null 2>&1; then
diff --git a/test/tileop_api/src/AssemblyLifetime.cpp b/test/tileop_api/src/AssemblyLifetime.cpp
deleted file mode 100644
index 433c7a2..0000000
--- a/test/tileop_api/src/AssemblyLifetime.cpp
+++ /dev/null
@@ -1,45 +0,0 @@
-// Object-level lifetime coverage. An OPEN assembly parent may not escape
-// through a Tile reference or ordinary Tile store before LAST.
-#include <common/pto_tileop.hpp>
-
-using namespace pto;
-using Part = Tile<Location::Vec, float, 4, 8, BLayout::RowMajor>;
-using Parent = Tile<Location::Vec, float, 8, 8, BLayout::RowMajor>;
-using GM = global_tensor<float, RowMajor<4, 8>>;
-
-#if defined(TEST_OPEN_ESCAPE)
-__attribute__((noinline)) void open_escape(GM &src, Parent &out) {
-  auto init = range::assemble(out);
-  TLOAD(init, src);
-}
-#elif defined(TEST_OPEN_STORE)
-__attribute__((noinline)) void open_store_before_last(GM &src, float *dst) {
-  Parent parent;
-  auto init = range::assemble(parent);
-  TLOAD(init, src);
-  global_tensor<float, RowMajor<8, 8>> output(dst);
-  TSTORE(output, parent);
-  auto last = range::assemble_last(parent);
-  TLOAD_ASS(last, src);
-}
-#else
-__attribute__((noinline)) void closed_load_session(GM &src, Parent &out) {
-  Parent parent;
-  auto init = range::assemble(parent);
-  TLOAD(init, src);
-  auto last = range::assemble_last(parent);
-  TLOAD_ASS(last, src);
-  TMOV(out, parent);
-}
-
-__attribute__((noinline)) void closed_tadd_session(Part &a, Part &b, Part &out) {
-  Part parent;
-  auto init = range::assemble(parent);
-  TADD(init, a, b);
-  auto middle = range::assemble_middle(parent);
-  TADD_ASS(middle, a, b);
-  auto last = range::assemble_last(parent);
-  TADD_ASS(last, a, b);
-  TMOV(out, parent);
-}
-#endif
diff --git a/test/tileop_api/src/DirectCubeLayout.cpp b/test/tileop_api/src/DirectCubeLayout.cpp
index b8e9984..6b68a75 100644
--- a/test/tileop_api/src/DirectCubeLayout.cpp
+++ b/test/tileop_api/src/DirectCubeLayout.cpp
@@ -90,7 +90,6 @@ __attribute__((noinline)) void gmov_cube_m32(CubeM32Vec &d, CubeM32Vec &s) {
 
 // --- 3a. reduce/expand (SFU) selects the operand layout too ---
 using Red32 = VecTileM32<float, 32, 1>;
-using ArgRed32 = VecTileM32<uint32_t, 32, 1>;
 using Red16 = VecTileM16<float, 16, 1>;
 using RedRowMajor = Tile<Location::Vec, float, 32, 1, BLayout::RowMajor>;
 using RedCol32 = VecTileM32<float, 1, 32>;   // 1 x N for TCOL*
@@ -103,7 +102,7 @@ __attribute__((noinline)) void rowsum_cube_m16(Red16 &d, CubeM16Vec &s) {
   TROWSUM(d, s);
 }
 
-__attribute__((noinline)) void rowargmax_cube_m32(ArgRed32 &d, CubeM32Vec &s) {
+__attribute__((noinline)) void rowargmax_cube_m32(Red32 &d, CubeM32Vec &s) {
   TROWARGMAX(d, s);
 }
 
diff --git a/test/tileop_api/src/RowArgNegatives.cpp b/test/tileop_api/src/RowArgNegatives.cpp
deleted file mode 100644
index f326d77..0000000
--- a/test/tileop_api/src/RowArgNegatives.cpp
+++ /dev/null
@@ -1,18 +0,0 @@
-#include <common/pto_tileop.hpp>
-
-using namespace pto;
-
-using Source = Tile<Location::Vec, __half, 16, 16,
-                    BLayout::RowMajor, 16, 16>;
-using BadDestination = Tile<Location::Vec, __half, 16, 1,
-                            BLayout::RowMajor, 16, 1>;
-
-#if defined(SHOULD_FAIL_TROWARGMAX_NON_U32)
-void bad_max(BadDestination &dst, Source &src) { TROWARGMAX(dst, src); }
-#endif
-
-#if defined(SHOULD_FAIL_TROWARGMIN_NON_U32)
-void bad_min(BadDestination &dst, Source &src) { TROWARGMIN(dst, src); }
-#endif
-
-int main() { return 0; }
diff --git a/test/tileop_api/src/RuntimePhysicalCols.cpp b/test/tileop_api/src/RuntimePhysicalCols.cpp
deleted file mode 100644
index 8af8c60..0000000
--- a/test/tileop_api/src/RuntimePhysicalCols.cpp
+++ /dev/null
@@ -1,36 +0,0 @@
-#include <common/pto_tileop.hpp>
-
-using Wide = Tile<Location::Vec, float, 32, 256,
-                  BLayout::CubeM32, DYNAMIC, DYNAMIC>;
-using Row = Tile<Location::Vec, float, 32, 1,
-                 BLayout::CubeM32, DYNAMIC, 1>;
-
-static_assert(pto_elementwise_user_dims::user_dims_legal<Wide>(6, 2, 6));
-static_assert(pto_elementwise_user_dims::user_dims_legal<Wide>(5, 2, 5));
-static_assert(!pto_elementwise_user_dims::user_dims_legal<Wide>(0, 2, 6));
-static_assert(!pto_elementwise_user_dims::user_dims_legal<Wide>(7, 2, 6));
-static_assert(!pto_elementwise_user_dims::user_dims_legal<Wide>(6, 2, 257));
-static_assert(!pto_elementwise_user_dims::user_dims_legal<Wide>(6, 33, 6));
-
-__attribute__((noinline)) void runtime_physical_cols(
-    Wide &dst, Wide &left, Wide &right, Row &row,
-    unsigned validCol, unsigned validRow, unsigned physicalCol) {
-  TEXPANDS(dst, 0.0f, validCol, validRow, physicalCol);
-  TADD(dst, left, right, validCol, validRow, physicalCol);
-  TSUB(dst, left, right, validCol, validRow, physicalCol);
-  TROWEXPANDMUL(dst, left, row, validCol, validRow, physicalCol);
-}
-
-__attribute__((noinline)) void odd_runtime_physical_cols(
-    Wide &dst, Wide &left, Wide &right) {
-  TADD(dst, left, right, 5, 2, 5);
-}
-
-// Zero overrides use the instance's valid dimensions for dynamic Tile types.
-__attribute__((noinline)) void runtime_valid_shape_defaults(
-    Wide &dst, Wide &left, Wide &right) {
-  TEXPANDS(dst, 0.0f, 0, 0, 6);
-  TADD(dst, left, right, 0, 0, 6);
-}
-
-int main() { return 0; }
diff --git a/test/tileop_api/src/TileArrayParticipantMask.cpp b/test/tileop_api/src/TileArrayParticipantMask.cpp
deleted file mode 100644
index f0af479..0000000
--- a/test/tileop_api/src/TileArrayParticipantMask.cpp
+++ /dev/null
@@ -1,48 +0,0 @@
-#include <common/pto_tileop.hpp>
-
-using namespace pto;
-
-using Leaf = ElementTile<uint32_t, 32>;
-using Parent = ElementTile<uint32_t, 128>;
-
-#ifndef PTO_TILEOP_API_HAS_TILE_ARRAY_PARTICIPANT_MASK
-#error "TileArray participant-mask capability gate is required"
-#endif
-
-static_assert(TileArray<Leaf, 1, 4>::participant_mask == 15);
-static_assert(TileArray<Leaf, 1, 4, 8>::participant_mask == 8);
-
-PTO_PE_DOMAIN(8)
-extern "C" void assemble_pe0(const uint32_t *input, uint32_t *output) {
-  TileArray<Leaf, 1, 4, 8> leaves;
-  for (unsigned part = 0; part < 4; ++part) {
-    Leaf leaf;
-    global_tensor<uint32_t, RowMajor<32, 4>> source(input + part);
-    TLOAD(leaf, source);
-    TMOV(leaves[0][part], leaf);
-  }
-  Parent assembled = TASSEMBLY<Parent>(std::move(leaves));
-  TSTORE(output, assembled, 128);
-  auto parts = TPARTVIEW<32>(assembled, 128);
-  Leaf selected;
-  auto first = parts.part(0);
-  TMOV(selected, first);
-  auto &elements = TPARTELEMENT(selected);
-#pragma pto element for
-  for (unsigned element = 0; element < 32; ++element)
-    elements[element] ^= 1u;
-  global_tensor<uint32_t, RowMajor<32, 1>> selected_output(output + 128);
-  TSTORE(selected_output, selected);
-}
-
-extern "C" void assemble_default(const uint32_t *input, uint32_t *output) {
-  TileArray<Leaf, 1, 4> leaves;
-  for (unsigned part = 0; part < 4; ++part) {
-    Leaf leaf;
-    global_tensor<uint32_t, RowMajor<32, 4>> source(input + part);
-    TLOAD(leaf, source);
-    TMOV(leaves[0][part], leaf);
-  }
-  Parent assembled = TASSEMBLY<Parent>(std::move(leaves));
-  TSTORE(output, assembled, 128);
-}
diff --git a/test/tileop_api/src/TileArrayWriterSize.cpp b/test/tileop_api/src/TileArrayWriterSize.cpp
index dd248ec..a96d8fc 100644
--- a/test/tileop_api/src/TileArrayWriterSize.cpp
+++ b/test/tileop_api/src/TileArrayWriterSize.cpp
@@ -10,8 +10,7 @@ using Fragment = global_tensor<float, RowMajor<32, 4>>;
 
 __attribute__((noinline)) void assemble_writer_size(
     Fragment &gm0, Fragment &gm1, Fragment &gm2, Fragment &gm3,
-    Parent &out) {
-  Parent parent;
+    Parent &parent) {
   auto first = range::assemble<4, 0>(parent);
   TLOAD(first, gm0);
 
@@ -23,7 +22,6 @@ __attribute__((noinline)) void assemble_writer_size(
 
   auto last = range::assemble_last<4, 12>(parent);
   TLOAD_ASS(last, gm3);
-  out = parent; // Publish the complete parent only after LAST.
 }
 
 __attribute__((noinline)) void assemble_single_writer(
diff --git a/test/tileop_api/src/TileRegionTCVTSubviewAssembly.cpp b/test/tileop_api/src/TileRegionTCVTSubviewAssembly.cpp
index f69a7d4..05d158d 100644
--- a/test/tileop_api/src/TileRegionTCVTSubviewAssembly.cpp
+++ b/test/tileop_api/src/TileRegionTCVTSubviewAssembly.cpp
@@ -3,8 +3,8 @@
 using namespace pto;
 using SourceParent = CubeTileM32<float, 32, 64>;
 using SourceFragment = CubeTileM32<float, 32, 16>;
-using DestinationFragment = CubeTileM32<__bf16, 32, 16>;
-using DestinationParent = CubeTileM32<__bf16, 32, 64>;
+using DestinationFragment = Tile<Location::Vec, __bf16, 32, 16, BLayout::RowMajor>;
+using DestinationParent = Tile<Location::Vec, __bf16, 32, 64, BLayout::RowMajor>;
 __attribute__((noinline)) DestinationParent convert_subviews(SourceParent &parent) {
   auto views = TPARTVIEW<SourceFragment, 1, 4>(parent);
   TileArray<DestinationFragment, 1, 4> output;
diff --git a/test/tileop_api/src/ValidShapePerDimSweep.cpp b/test/tileop_api/src/ValidShapePerDimSweep.cpp
index 347d5ef..43ed454 100644
--- a/test/tileop_api/src/ValidShapePerDimSweep.cpp
+++ b/test/tileop_api/src/ValidShapePerDimSweep.cpp
@@ -170,8 +170,6 @@ template <int VR,int VC> using T =
   Tile<Location::Vec, float, 32, 32, BLayout::RowMajor, VR, VC>;
 template <int VR,int VC> using C1 =
   Tile<Location::Vec, float, 32, 1, BLayout::RowMajor, VR, VC>;
-template <int VR,int VC> using U1 =
-  Tile<Location::Vec, uint32_t, 32, 1, BLayout::RowMajor, VR, VC>;
 
 // EXPAND (dst=T, src=C1 单列)
 #define SWEX(FN) \
@@ -190,10 +188,10 @@ SWEX3(TROWEXPANDMAX) SWEX3(TROWEXPANDMIN) SWEX3(TROWEXPANDEXPDIF)
 
 // ARGMAX/ARGMIN (dst=C1, src=T)
 #define SWAR(FN) \
-  __attribute__((noinline)) void FN##_ss(U1<16,1> &d,T<16,32> &s){FN(d,s);} \
-  __attribute__((noinline)) void FN##_sd(U1<16,-1> &d,T<16,-1> &s){FN(d,s);} \
-  __attribute__((noinline)) void FN##_ds(U1<-1,1> &d,T<-1,32> &s){FN(d,s);} \
-  __attribute__((noinline)) void FN##_dd(U1<-1,-1> &d,T<-1,-1> &s){FN(d,s);}
+  __attribute__((noinline)) void FN##_ss(C1<16,1> &d,T<16,32> &s){FN(d,s);} \
+  __attribute__((noinline)) void FN##_sd(C1<16,-1> &d,T<16,-1> &s){FN(d,s);} \
+  __attribute__((noinline)) void FN##_ds(C1<-1,1> &d,T<-1,32> &s){FN(d,s);} \
+  __attribute__((noinline)) void FN##_dd(C1<-1,-1> &d,T<-1,-1> &s){FN(d,s);}
 SWAR(TROWARGMAX) SWAR(TROWARGMIN)
 
 
diff --git a/test/tileop_api/verify_target_cxx_frontend.sh b/test/tileop_api/verify_target_cxx_frontend.sh
index 52e6517..346ec72 100755
--- a/test/tileop_api/verify_target_cxx_frontend.sh
+++ b/test/tileop_api/verify_target_cxx_frontend.sh
@@ -12,8 +12,7 @@ trap 'rm -rf "$OUT"' EXIT
 # Use a copied resource tree whose tileop-api subtree is this checkout, so the
 # verification cannot accidentally exercise stale installed headers.
 RESOURCE_DIR="$OUT/clang-resource"
-BASE_RESOURCE_DIR=${CLANG_RESOURCE_DIR:-$("$TC_DIR/clang++" -print-resource-dir)}
-cp -R "$BASE_RESOURCE_DIR" "$RESOURCE_DIR"
+cp -R "$("$TC_DIR/clang++" -print-resource-dir)" "$RESOURCE_DIR"
 rm -rf "$RESOURCE_DIR/include/tileop-api"
 cp -R "$ROOT/include" "$RESOURCE_DIR/include/tileop-api"
 
@@ -53,8 +52,7 @@ for source in RangeSubview.cpp GMov.cpp TileRegionUnary.cpp \
               Issue288CubeM32BinaryAssembly.cpp \
               TOrAssSubview.cpp \
               TileArrayTCVTE8M0.cpp \
-              SharedTransposeNonSquare.cpp TCI.cpp TRowExpandMul.cpp \
-              RuntimePhysicalCols.cpp; do
+              SharedTransposeNonSquare.cpp TCI.cpp TRowExpandMul.cpp; do
   "$TC_DIR/clang++" "${FLAGS[@]}" -fsyntax-only \
     "$ROOT/test/tileop_api/src/$source"
 done
@@ -216,155 +214,6 @@ for op, offset, last in zip(ops, (4, 12), (0, 1)):
         raise SystemExit("TileArrayWriterSize: Shared continuation rebound its parent")
 PY_ASSEMBLE
 
-"$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-  "$ROOT/test/tileop_api/src/RuntimePhysicalCols.cpp" \
-  -o "$OUT/RuntimePhysicalCols.ll"
-python3 - "$OUT/RuntimePhysicalCols.ll" <<'PY_RUNTIME_COLS'
-import re
-import sys
-from pathlib import Path
-
-text = Path(sys.argv[1]).read_text(encoding="utf-8")
-body = next(part for part in re.split(r"(?=^define )", text, flags=re.MULTILINE)
-            if part.startswith("define ") and "runtime_physical_cols" in part)
-operations = [part for part in re.split(r"(?=BSTART\.TEPL)", body)
-              if "BSTART.TEPL" in part]
-runtime_operations = [operation for operation in operations
-                      if "B.DIM $" in operation and ", 0, ->lb2" in operation]
-if len(runtime_operations) != 4:
-    raise SystemExit(
-        f"expected four register-LB2 operations, got {len(runtime_operations)}")
-if sum("B.IOR" in operation for operation in runtime_operations) != 1:
-    raise SystemExit("runtime-geometry fixture lost scalar-fill binding")
-odd = next(part for part in re.split(r"(?=^define )", text, flags=re.MULTILINE)
-           if part.startswith("define ") and "odd_runtime_physical_cols" in part)
-if "BSTART.TEPL" not in odd or "->lb2" not in odd:
-    raise SystemExit("odd runtime PhysicalCol fixture was not emitted")
-defaults = next(part for part in re.split(r"(?=^define )", text, flags=re.MULTILINE)
-                if part.startswith("define ") and "runtime_valid_shape_defaults" in part)
-if defaults.count("BSTART.TEPL") != 2:
-    raise SystemExit("dynamic zero overrides must use instance dimensions, not DYNAMIC=-1")
-PY_RUNTIME_COLS
-
-for source in TileRegionUnary TileRegionCubeSubview Issue166RmsNormB128; do
-  "$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-    "$ROOT/test/tileop_api/src/$source.cpp" -o "$OUT/$source.ll"
-  "$TC_DIR/clang++" "${FLAGS[@]}" -c \
-    "$ROOT/test/tileop_api/src/$source.cpp" -o "$OUT/$source.o"
-done
-"$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-  "$ROOT/test/tileop_api/src/RangeAssemble.cpp" -o "$OUT/RangeAssemble.ll"
-"$TC_DIR/clang++" "${FLAGS[@]}" -c \
-  "$ROOT/test/tileop_api/src/AssemblyLifetime.cpp" -o "$OUT/AssemblyLifetime.o"
-# RangeAssemble contains isolated INIT/MIDDLE syntax probes; only complete
-# sessions are legal object-level programs. Keep escaping OPEN parents negative.
-for mode in TEST_OPEN_ESCAPE TEST_OPEN_STORE; do
-  if "$TC_DIR/clang++" "${FLAGS[@]}" -D"$mode" -c \
-      "$ROOT/test/tileop_api/src/AssemblyLifetime.cpp" \
-      -o "$OUT/$mode.o" >"$OUT/$mode.log" 2>&1; then
-    echo "FAIL: $mode accepted an OPEN assembly parent escape" >&2
-    exit 1
-  fi
-  grep -Fq "cannot use an OPEN B.ASSEMBLE parent" "$OUT/$mode.log"
-done
-python3 - "$OUT/TileRegionUnary.ll" "$OUT/TileRegionCubeSubview.ll" \
-          "$OUT/Issue166RmsNormB128.ll" "$OUT/RangeAssemble.ll" <<'PY'
-import re
-import sys
-from pathlib import Path
-
-
-def region_operations(path):
-    return [line for line in Path(path).read_text(encoding="utf-8").splitlines()
-            if "asm sideeffect" in line and "BSTART.TEPL" in line
-            and "B.SUBVIEW" in line]
-
-
-unary = region_operations(sys.argv[1])
-if len(unary) != 9:
-    raise SystemExit(f"TileRegionUnary: expected 9 region operations, got {len(unary)}")
-for operation in unary:
-    if "i32 16, i32 32, i32 16, i32 5, i32 5" not in operation:
-        raise SystemExit("TileRegionUnary: LB2 does not use the 16-column fragment")
-
-cube = region_operations(sys.argv[2])
-if len(cube) != 1:
-    raise SystemExit(f"TileRegionCubeSubview: expected 1 region operation, got {len(cube)}")
-if "i32 16, i32 32, i32 16, i32 5, i32 5" not in cube[0]:
-    raise SystemExit("TileRegionCubeSubview: LB2 does not use the 16-column fragment")
-if "i32 16, i32 32, i32 64, i32 5, i32 5" in cube[0]:
-    raise SystemExit("TileRegionCubeSubview: LB2 leaked the 64-column parent geometry")
-
-rowsum = [line for line in region_operations(sys.argv[3]) if "i32 64, i64" in line]
-expected = (
-    "i32 8, i32 32, i32 8, i32 4, i32 4, i32 64, i64 0",
-    "i32 8, i32 8, i32 8, i32 2, i32 2, i32 64, i64 0",
-)
-if len(rowsum) != len(expected):
-    raise SystemExit(f"Issue166RmsNormB128: expected 2 TROWSUM regions, got {len(rowsum)}")
-for operands in expected:
-    if not any(operands in operation for operation in rowsum):
-        raise SystemExit(f"Issue166RmsNormB128: missing source geometry {operands}")
-
-range_assemble = Path(sys.argv[4]).read_text(encoding="utf-8")
-body = next((part for part in re.split(r"(?=^define )", range_assemble,
-                                       flags=re.MULTILINE)
-             if "assemble_tadd_session" in part), None)
-if body is None:
-    raise SystemExit("RangeAssemble: missing TADD assembly-session fixture")
-operations = [line for line in body.splitlines()
-              if "asm sideeffect" in line and "BSTART.TEPL" in line]
-if len(operations) != 3:
-    raise SystemExit(f"RangeAssemble: expected 3 TADD operations, got {len(operations)}")
-if "B.IOT $5.reuse, $6.reuse, mask=1111, last, ->$0" not in operations[0]:
-    raise SystemExit("RangeAssemble: INIT TADD must allocate its destination")
-if "B.ASSEMBLE ${8:c}, ${9:c}" not in operations[0]:
-    raise SystemExit("RangeAssemble: INIT TADD is missing B.ASSEMBLE")
-for operation in operations[1:]:
-    if "B.IOT $4.reuse, $5.reuse, mask=1111" not in operation or \
-       "B.IOT $7.reuse, mask=1111, last" not in operation:
-        raise SystemExit("RangeAssemble: continuation TADD has invalid binders")
-    if "B.ASSEMBLE" not in operation:
-        raise SystemExit("RangeAssemble: continuation TADD is missing B.ASSEMBLE")
-PY
-
-# Typed factories carry a seventh WriterSizeCode template argument. Verify
-# continuation overload selection and real object code for external parents.
-"$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-  "$ROOT/test/tileop_api/src/TileArrayWriterSize.cpp" -o "$OUT/TileArrayWriterSize.ll"
-"$TC_DIR/clang++" "${FLAGS[@]}" -c \
-  "$ROOT/test/tileop_api/src/TileArrayWriterSize.cpp" -o "$OUT/TileArrayWriterSize.o"
-# Shared references still have a separate backend handle-copy gap; verify
-# descriptor lowering in IR without claiming Shared object/runtime closure.
-"$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-  "$ROOT/test/tileop_api/src/SharedTloadAssembly.cpp" -o "$OUT/SharedTloadAssembly.ll"
-python3 - "$OUT/TileArrayWriterSize.ll" "$OUT/SharedTloadAssembly.ll" <<'PY_ASSEMBLE'
-import re
-import sys
-from pathlib import Path
-text = Path(sys.argv[1]).read_text()
-body = next(part for part in re.split(r"(?=^define )", text, flags=re.MULTILINE)
-            if part.startswith("define ") and "assemble_writer_size" in part)
-ops = [line for line in body.splitlines() if "asm sideeffect" in line]
-if len(ops) != 4 or any("B.ASSEMBLE" not in op for op in ops):
-    raise SystemExit("TileArrayWriterSize: an INIT/MIDDLE/LAST descriptor was lost")
-for op, offset, last in zip(ops[1:], (4, 8, 12), (0, 0, 1)):
-    if "->" in op.split("B.IOT", 1)[1].split("B.ASSEMBLE", 1)[0]:
-        raise SystemExit("TileArrayWriterSize: continuation allocated a new parent")
-    if f"i32 0, i32 {last}, i32 {offset}, i32 3)" not in op:
-        raise SystemExit("TileArrayWriterSize: incorrect continuation extent or state")
-text = Path(sys.argv[2]).read_text()
-body = next(part for part in re.split(r"(?=^define )", text, flags=re.MULTILINE)
-            if part.startswith("define ") and "assemble_shared_continuations" in part)
-ops = [line for line in body.splitlines() if "asm sideeffect" in line]
-if len(ops) != 2 or any("B.ASSEMBLE" not in op or "B.IOS" not in op for op in ops):
-    raise SystemExit("TileArrayWriterSize: Shared continuation lost its descriptor")
-for op, offset, last in zip(ops, (4, 12), (0, 1)):
-    if "->" in op.split("B.IOS", 1)[1].split("B.ASSEMBLE", 1)[0] or \
-       f"i32 0, i32 {last}, i32 {offset}, i32 3)" not in op:
-        raise SystemExit("TileArrayWriterSize: Shared continuation rebound its parent")
-PY_ASSEMBLE
-
 "$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
   "$ROOT/test/tileop_api/src/TRowExpandMul.cpp" -o "$OUT/TRowExpandMul.ll"
 python3 - "$OUT/TRowExpandMul.ll" <<'PY'
@@ -563,16 +412,12 @@ for name in ("tor_ass_subview", "tor_ass_explicit_range"):
         raise SystemExit(f"{name}: invalid source/destination modifier order")
 PY
 
-"$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
-  "$ROOT/test/tileop_api/src/TileRegionTCVTSubviewAssembly.cpp" \
-  -o "$OUT/TileRegionTCVTSubviewAssembly.ll"
 "$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
   "$ROOT/test/tileop_api/src/TileRegionTCVTSubview.cpp" \
   -o "$OUT/TileRegionTCVTSubview.ll"
 "$TC_DIR/clang++" "${FLAGS[@]}" -S -emit-llvm \
   "$ROOT/test/tileop_api/src/TileRegionShift.cpp" -o "$OUT/TileRegionShift.ll"
-python3 - "$OUT/TileRegionTCVTSubview.ll" \
-  "$OUT/TileRegionTCVTSubviewAssembly.ll" "$OUT/TileRegionShift.ll" <<'PY'
+python3 - "$OUT/TileRegionTCVTSubview.ll" "$OUT/TileRegionShift.ll" <<'PY'
 import re
 import sys
 from pathlib import Path
@@ -585,19 +430,7 @@ if "BSTART.TEPL 27" not in body or body.count("B.SUBVIEW") != 1:
 if "B.ASSEMBLE" in body or "B.DIM zero, $2, ->lb2" in body:
     raise SystemExit("convert_subview: ordinary CUBE TCVT emitted assembly/LB2")
 
-assembled_tcvt = Path(sys.argv[2]).read_text(encoding="utf-8")
-assembled_body = next(
-    part for part in re.split(r"(?=^define )", assembled_tcvt,
-                              flags=re.MULTILINE)
-    if "convert_subviews" in part)
-if assembled_body.count("B.SUBVIEW") != 4 or assembled_body.count("B.ASSEMBLE") != 4:
-    raise SystemExit("convert_subviews: invalid TCVT modifier counts")
-for line in (line for line in assembled_body.splitlines()
-             if "asm sideeffect" in line and "BSTART.TEPL 27" in line):
-    if "->lb2" in line:
-        raise SystemExit("convert_subviews: CUBE TCVT emitted forbidden LB2")
-
-shift = Path(sys.argv[3]).read_text(encoding="utf-8")
+shift = Path(sys.argv[2]).read_text(encoding="utf-8")
 functions = re.split(r"(?=^define )", shift, flags=re.MULTILINE)
 ordinary = next(part for part in functions if "shift_subview" in part and
                 "assemble_shift_subviews" not in part)
@@ -678,6 +511,6 @@ PY
 grep -Eiq 'BSTART\.CUBE[[:space:]]+TMATMULMX,[[:space:]]+FP16' "$OUT/MXScaleVariants.dis"
 grep -Eiq 'BSTART\.CUBE[[:space:]]+TMATMULMX\.ACC,[[:space:]]+E4M3' "$OUT/MXScaleVariants.dis"
 grep -Eiq 'BSTART\.CUBE[[:space:]]+TGEMVMX\.BIAS,[[:space:]]+BF16' "$OUT/MXScaleVariants.dis"
-grep -Eq 'B\.IOT[[:space:]]+t#[1-8]\.reuse, t#[1-8]\.reuse, mask=1111' "$OUT/MXScaleVariants.dis"
+grep -Eq 'B\.IOT[[:space:]]+t#[1-8], t#[1-8], mask=1111' "$OUT/MXScaleVariants.dis"
 
 echo "Linx target C++ frontend MX/effective-shape contract: PASS"