TileOP: enforce CUBE InternalAcc CCTRL legality (spec#236, Phase 1)

include/common/pto_tile.hpp

@@ -156,8 +156,10 @@ struct FixpAttr {
   // (raw accumulator D) forbids final post-process; CCTRL[1] (ACC-only
   // transparent cache hint) is validated against the operation kind at the
   // wrapper layer.
-  static constexpr FixpAttr with_cube_ctrl(CubeControl Ctrl) {
-    FixpAttr Attr;
+  // Chainable: preserves all other attributes, unlike the parameter-free
+  // static constructors (f16/bf16/keep_acc) which start from defaults.
+  constexpr FixpAttr with_cube_ctrl(CubeControl Ctrl) const {
+    FixpAttr Attr = *this;
     Attr.CubeCtrl = Ctrl;
     return Attr;
   }

include/jcore/template_asm.hpp

@@ -3944,6 +3944,7 @@ PTO_SHARED_INLINE void matmul(Dst &dst, A &a, B &b, size_t M, size_t N,
   // Local-A/Shared-B the A shard descriptor is per-PE (valid_rows == pe_m)
   // and cannot supply group_M, so the caller must pass it here.
   validate_matrix_contract<Attr, Dst, A, B>();
+  validate_cube_ctrl_contract<Attr, /*IsAccForm=*/false>();
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {
     asm volatile(
         PTO_MATMUL_HEADER("TMATMUL", PTO_FIXP_ATTR)
@@ -4058,6 +4059,7 @@ PTO_SHARED_INLINE void matmul_acc(Dst &dst, C &c, A &a, B &b, size_t M,
                                   size_t N, size_t K) {
   validate_matrix_contract<Attr, Dst, A, B>();
   validate_matrix_accumulator_contract<Attr, Dst, C, A, B>();
+  validate_cube_ctrl_contract<Attr, /*IsAccForm=*/true>();
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {
     asm volatile(
         PTO_MATMUL_HEADER("TMATMUL.ACC", PTO_FIXP_ATTR)
@@ -5590,7 +5592,7 @@ PTO_SHARED_INLINE void emit_fixp(
     uint64_t quant_gpr, uint64_t lrelu_gpr, size_t M, size_t N, size_t K) {
   validate_matrix_contract<Attr, Dst, A, B>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, A, B,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, false>();
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {
     PTO_FIXP_DISPATCH(PTO_FIXP_EMIT_LOCAL);
   } else if constexpr (is_shared_tile_v<A> && !is_shared_tile_v<B>) {
@@ -5625,7 +5627,7 @@ PTO_SHARED_INLINE void emit_matmul_acc_fixp(
   validate_matrix_accumulator_contract<Attr, Dst, C_, A, B>();
   validate_cscale_contract<Attr, C_, CScale>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, A, B,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, true>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, true>();
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {
     PTO_FIXP_DISPATCH(PTO_FIXP_ACC_EMIT_LOCAL);
   } else if constexpr (is_shared_tile_v<A> && !is_shared_tile_v<B>) {
@@ -5649,7 +5651,7 @@ PTO_SHARED_INLINE void emit_matmul_bias_fixp(
   validate_matrix_contract<Attr, Dst, A, B>();
   validate_matrix_bias_contract<Attr, BiasT, A, B>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, A, B,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, false>();
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {
     PTO_FIXP_DISPATCH(PTO_FIXP_BIAS_EMIT_LOCAL);
   } else if constexpr (is_shared_tile_v<A> && !is_shared_tile_v<B>) {
@@ -5839,7 +5841,7 @@ PTO_SHARED_INLINE void emit_gemv_fixp(
   uint64_t quant_gpr, uint64_t lrelu_gpr, size_t M, size_t N, size_t K) {
   validate_gemv_contract<Attr, Dst, Vec, Mtx>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, Vec, Mtx,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, false>();
   PTO_FIXP_DISPATCH(PTO_FIXP_GV_GV_EMIT_LOCAL);
 }
 
@@ -5858,7 +5860,7 @@ PTO_SHARED_INLINE void emit_gemv_bias_fixp(
   validate_gemv_contract<Attr, Dst, Vec, Mtx>();
   validate_matrix_bias_contract<Attr, BiasT, Vec, Mtx>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, Vec, Mtx,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, false>();
   PTO_FIXP_DISPATCH(PTO_FIXP_GV_GVB_EMIT_LOCAL);
 }
 
@@ -5877,7 +5879,7 @@ PTO_SHARED_INLINE void emit_gemv_acc_fixp(
   validate_gemv_contract<Attr, Dst, Vec, Mtx>();
   validate_matrix_accumulator_contract<Attr, Dst, C, Vec, Mtx>();
   validate_matrix_postprocess_contract<Attr, SrcMask, OutMask, Vec, Mtx,
-      RowIn, QuantTile, ReluTile, RowOut, GroupOut, true>();
+      RowIn, QuantTile, ReluTile, RowOut, GroupOut, false, true>();
   PTO_FIXP_DISPATCH(PTO_FIXP_GV_GVA_EMIT_LOCAL);
 }
 
@@ -6025,6 +6027,7 @@ PTO_SHARED_INLINE void Name(Dst &dst, A &a, ScaleA &scale_a, B &b,             \
   constexpr bool HasScaleB = true;                                               \
   validate_matrix_contract<Attr, Dst, A, B, true>();                                      \
   validate_matrix_scale_contract<Attr, HasScaleA, HasScaleB, ScaleA, A, ScaleB, B>();      \
+  validate_cube_ctrl_contract<Attr, /*IsAccForm=*/false>();                     \
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {                \
     asm volatile(                                                               \
         PTO_MATMUL_HEADER(Opcode, PTO_FIXP_ATTR)                               \
@@ -6145,6 +6148,7 @@ PTO_SHARED_INLINE void Name(Dst &dst, A &a, ScaleA &scale_a, B &b,             \
     validate_matrix_accumulator_contract<Attr, Dst, Extra, A, B, true>();                       \
   else                                                                                    \
     validate_matrix_bias_contract<Attr, Extra, A, B, true>();                                   \
+  validate_cube_ctrl_contract<Attr, IsAcc>();                                   \
   if constexpr (!is_shared_tile_v<A> && !is_shared_tile_v<B>) {                \
     asm volatile(                                                               \
         PTO_MATMUL_HEADER(Opcode, PTO_FIXP_ATTR)                               \

test/tileop_api/compile.all

@@ -25,7 +25,7 @@ ACTIVE_FIXTURES=(
   TAdd_mask TAdd TAdds TAnd TLoad TStore TCvt TDiv TDivs TExp TMax TMaxs
   TMov TMul TMuls TOr TRecip TRem TRowMax TRowSum TSqrt TSub TSubs TTrans
   RangeSubview RangeAssemble SharedRange TileRegion TileArray TileArrayCube TileArrayAssemblyOffsets TileArrayRegionAsm TileRegionUnaryAssembly TileRegionUnary TileRegionBinary TileRegionBinaryAssembly TileRegionUnarySubviewAssembly TileRegionScalarAssembly TileRegionTCVTSubviewAssembly TileRegionSharedSubview TileRegionCubeSubview
-  ValidShapeImmediate ValidShapePerDim ValidShapePerDimSweep TSELCanonical
+  ValidShapeImmediate ValidShapePerDim ValidShapePerDimSweep CubeInternalAcc TSELCanonical
 )
 
 # Pre-v0.58 wrapper designs remain available to cpu_sim only and are not part

test/tileop_api/src/CubeInternalAcc.cpp

@@ -0,0 +1,35 @@
+// CUBE InternalAcc (PTO-ISA 0.58.6 spec#236) regression: CCTRL encoding and
+// legality. Positive cases lower the four CCTRL values into the B.DATR
+// PadValueOrByteId[1:0] union field; illegal combinations must be rejected
+// at compile time (see the paired negative runner).
+#include <jcore/template_asm.hpp>
+
+using namespace pto;
+
+using A = SharedMatrixLeft<float, 64, 16>;
+using B = SharedMatrixRight<float, 16, 16>;
+using C = CubeAccumulatorM16<float, 16, 16>;
+using Opt = fixp::Options<FixpAttr{}>;
+
+// CCTRL = 0 (None)    -> B.DATR ..., Zero, ...
+__attribute__((noinline)) void acc_none(C &d, C &c, A &a, B &b) {
+  auto sa = TMOV_L2S_INSERT(a); auto sb = TMOV_L2S_INSERT(b);
+  TMATMUL_ACC(d, c, sa, sb, Opt{});
+}
+// CCTRL = 2 (InternalAccHint, ACC-only transparent cache hint) -> Min
+__attribute__((noinline)) void acc_hint(C &d, C &c, A &a, B &b) {
+  auto sa = TMOV_L2S_INSERT(a); auto sb = TMOV_L2S_INSERT(b);
+  TMATMUL_ACC(d, c, sa, sb, Opt{}.acc_hint());
+}
+// CCTRL = 1 (RawAccumulator, raw accumulator-type D) -> Max
+__attribute__((noinline)) void acc_raw(C &d, C &c, A &a, B &b) {
+  auto sa = TMOV_L2S_INSERT(a); auto sb = TMOV_L2S_INSERT(b);
+  TMATMUL_ACC(d, c, sa, sb, Opt{}.raw_acc());
+}
+// CCTRL = 3 (RawAccumulator | InternalAccHint) -> Null
+__attribute__((noinline)) void acc_raw_hint(C &d, C &c, A &a, B &b) {
+  auto sa = TMOV_L2S_INSERT(a); auto sb = TMOV_L2S_INSERT(b);
+  TMATMUL_ACC(d, c, sa, sb, Opt{}.raw_acc().acc_hint());
+}
+
+int main() { return 0; }

test/tileop_api/src/CubeInternalAccNegatives.cpp

@@ -0,0 +1,33 @@
+// CUBE InternalAcc negative contracts: each NEG_CASE must fail to compile.
+// Driven by run_negatives-style invocations (-DNEG_CASE=N).
+#include <jcore/template_asm.hpp>
+
+using namespace pto;
+
+using A = SharedMatrixLeft<float, 64, 16>;
+using B = SharedMatrixRight<float, 16, 16>;
+using C = CubeAccumulatorM16<float, 16, 16>;
+using CH = Tile<Location::Acc, __half, 16, 16, BLayout::CubeM16>;
+
+#ifdef NEG_CASE
+__attribute__((noinline)) void neg(
+#ifdef NEG_CASE2_DST
+    CH &d,
+#else
+    C &d,
+#endif
+    C &c, A &a, B &b) {
+  auto sa = TMOV_L2S_INSERT(a); auto sb = TMOV_L2S_INSERT(b);
+#if NEG_CASE == 1
+  // CCTRL[1] hint on a non-ACC TMATMUL: must be rejected.
+  TMATMUL(d, sa, sb, fixp::Options<FixpAttr{}.with_cube_ctrl(CubeCtrlInternalAccHint)>{});
+#elif NEG_CASE == 2
+  // CCTRL[0] raw D + f16 PreQuant (dtype-matching dst): must be rejected.
+  TMATMUL_ACC(d, c, sa, sb, fixp::Options<FixpAttr::f16().with_cube_ctrl(CubeCtrlRawAccumulator)>{});
+#elif NEG_CASE == 3
+  // CCTRL[0] raw D + Relu: must be rejected.
+  TMATMUL_ACC(d, c, sa, sb, fixp::Options<FixpAttr::keep_acc(FixpReluMode::Relu).with_cube_ctrl(CubeCtrlRawAccumulator)>{});
+#endif
+}
+int main() { return 0; }
+#endif