rocm/gfx1151: K-unroll WMMA GEMM lever — measured, regresses, kept as reference

docs/audit/backend/rocm/STRIX_HALO_EXECUTION_PLAN.md

@@ -288,6 +288,33 @@ lever has a cliff. Correctness of both tiles is guarded by
 `test_rocm_wmma_runtime_symbol.py` (small → 2×4, the 1024³ ragged case → 3×4).
 This supersedes the staging rungs (2/3) as the production perf story on this APU.
 
+**Stage H addendum — asymmetric-tile + K-unroll sweep (2026-07-10, negative).**
+Two remaining occupancy levers, swept on-device (best-of-N, kernel-only):
+
+- *Every untested asymmetric macro-tile* (2×3, 3×3, 4×2, 2×5, 5×2, 5×3, 3×5, 2×6,
+  6×2, 2×7, 2×8, 3×6) at 1024–4096³. **No untested tile beats the 2×4/3×4
+  adaptive** — 3×4 is the confirmed sweet spot 1024–3072³; 4×3 edges 3×4 by ~3%
+  at 4096³ only (a wash at 3072³). The register-budget sweet spot is genuinely 12
+  tiles.
+- *K-unrolling* — a new reference rung (`kKernelTemplateKU`, `..._bench_ku` /
+  `..._ku` C-ABI, verified correct by `test_rocm_gemm_ku_reference.py`): process
+  KU 16-wide K-panels per step (16·KU contiguous loads → more MLP) then KU·MT·NT
+  WMMAs. **Regresses.** KU2 loses −54…−66% on 3×4 and −43…−49% on 4×3; KU4
+  collapses −60…−71% everywhere; the only positive is 2×4-KU2 at 4096³ (+2%, still
+  far below 3×4-KU1). Root cause is the same VGPR/occupancy limit — the KU× a/b
+  load buffers blow the register budget and collapse waves/CU.
+
+**Synthesis:** every lever that adds register pressure — bigger tiles (4×4),
+K-unroll, register prefetch — hits the same occupancy cliff, and every
+memory-hierarchy lever (LDS, pipelined LDS) is a wash on unified LPDDR5x. The
+production register kernel is at the practical ceiling for this design on Strix
+Halo (~13 TFLOP/s ≈ 22% of the ~59 f16 WMMA peak). The named **dual-issue** lever
+does not apply: the inner loop is WMMA-issue-bound (12 back-to-back WMMAs/step),
+and RDNA3 dual-issue (VOPD) pairs *VALU* ops, not WMMA. Past here needs a
+fundamentally lower-register-footprint design (or a datacenter arch where global
+bandwidth is the bottleneck and the kept LDS/pipe/KU rungs pay off). The KU rung
+is retained as a correctness-verified reference for those targets.
+
 ### Stage G — flash_attn executes on gfx1151 (2026-06-23): second op after matmul
 
 `flash_attn` now executes natively on the AMD GPU — the **second op after

docs/audit/generated/compiler_progress.csv

@@ -16,7 +16,7 @@ primitive_axis,backend_kernel,mixed,99,480,381,"primitive contract axis; open me
 integration,Verifier coverage,mixed,163,174,11,"no_verifier=11, real=163",docs/audit/generated/verifier_coverage.csv,"Add real verifier implementations for no_verifier ops, prioritizing native codegen lanes."
 integration,Direct test evidence,mixed,363,480,117,"covered_by_family=26, directly_tested=363, hardware_gated=4, structural_only=87",docs/audit/generated/test_coverage.csv,Convert structural_only and needs_direct_test rows into direct compare fixtures; keep hardware_gated tied to backend proof.
 integration,Runtime execution matrix,closed,151,151,0,"apple_cpu=2, apple_gpu=16, cpu=2, nvidia_sm120=1, rocm=68, x86=62",docs/audit/generated/runtime_execution_matrix.csv,Add rows only when a launch path actually executes.
-integration,Runtime ABI symbols,mixed,397,650,253,"apple=557, nvidia=7, rocm=10, x86=76",docs/audit/generated/runtime_abi.csv,Reduce stub-only ABI rows where a backend claims native execution.
+integration,Runtime ABI symbols,mixed,399,652,253,"apple=557, nvidia=7, rocm=12, x86=76",docs/audit/generated/runtime_abi.csv,Reduce stub-only ABI rows where a backend claims native execution.
 integration,Audited repo surfaces,mixed,31,58,27,"archived=4, compile_only=12, runnable=31, runnable_optional=1, scaffold=10",docs/audit/generated/surface_status.csv,Graduate compile_only/scaffold entries that exercise compiler pathways; archive dead surfaces.
 codegen_pathway,Apple CPU,closed,174,174,0,"runtime: apple_cpu=2; target_map: fused=2, reference=170",docs/audit/generated/apple_target_map.csv,Keep as regression baseline for CPU value-call/runtime ABI.
 codegen_pathway,Apple GPU,mixed,187,188,1,"runtime: apple_gpu=16; target_map: absent=1, compiled=102, fused=58, hardware_verified=11",docs/audit/generated/apple_target_map.csv,Close the remaining absent target-map lane or document why it is host-only.

docs/audit/generated/compiler_progress.md

@@ -43,7 +43,7 @@ A row is not marked incomplete merely because Apple, x86, ROCm, and CUDA are not
 | `Verifier coverage` | mixed | 163 | 174 | 11 | no_verifier=11, real=163 | Add real verifier implementations for no_verifier ops, prioritizing native codegen lanes. |
 | `Direct test evidence` | mixed | 363 | 480 | 117 | covered_by_family=26, directly_tested=363, hardware_gated=4, structural_only=87 | Convert structural_only and needs_direct_test rows into direct compare fixtures; keep hardware_gated tied to backend proof. |
 | `Runtime execution matrix` | closed | 151 | 151 | 0 | apple_cpu=2, apple_gpu=16, cpu=2, nvidia_sm120=1, rocm=68, x86=62 | Add rows only when a launch path actually executes. |
-| `Runtime ABI symbols` | mixed | 397 | 650 | 253 | apple=557, nvidia=7, rocm=10, x86=76 | Reduce stub-only ABI rows where a backend claims native execution. |
+| `Runtime ABI symbols` | mixed | 399 | 652 | 253 | apple=557, nvidia=7, rocm=12, x86=76 | Reduce stub-only ABI rows where a backend claims native execution. |
 | `Audited repo surfaces` | mixed | 31 | 58 | 27 | archived=4, compile_only=12, runnable=31, runnable_optional=1, scaffold=10 | Graduate compile_only/scaffold entries that exercise compiler pathways; archive dead surfaces. |
 
 ## Code Generation Pathways

docs/audit/generated/runtime_abi.csv

@@ -568,9 +568,11 @@ rocm,tessera_rocm_wmma_flash_attn_f16,rocm_wmma_flash_attn,f16,src/compiler/code
 rocm,tessera_rocm_wmma_gemm_bf16,rocm_wmma_gemm,bf16,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16,rocm_wmma_gemm,f16,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_bench,rocm_wmma_gemm_f16_bench,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
+rocm,tessera_rocm_wmma_gemm_f16_bench_ku,rocm_wmma_gemm_f16_bench_ku,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_bench_lds,rocm_wmma_gemm_f16_bench_lds,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_bench_pipe,rocm_wmma_gemm_f16_bench_pipe,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_e2e_bench,rocm_wmma_gemm_f16_e2e_bench,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
+rocm,tessera_rocm_wmma_gemm_f16_ku,rocm_wmma_gemm_f16_ku,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_lds,rocm_wmma_gemm_f16_lds,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 rocm,tessera_rocm_wmma_gemm_f16_pipe,rocm_wmma_gemm_f16_pipe,,src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp
 x86,tessera_x86_amx_available,amx_available,,src/compiler/codegen/tessera_x86_backend/include/tessera/x86/target.h

docs/audit/generated/runtime_abi.md

@@ -4,7 +4,7 @@ Human-readable view. The canonical machine-readable artifact is `runtime_abi.csv
 
 ## Headline
 
-- **393** unique `extern "C" tessera_*` C ABI symbols across all backends.
+- **395** unique `extern "C" tessera_*` C ABI symbols across all backends.
 - **6 / 6** core runtime headers present.
 - **134** Apple GPU kernel families with per-dtype variants.
 
@@ -25,7 +25,7 @@ Human-readable view. The canonical machine-readable artifact is `runtime_abi.csv
 |---------|-------------------------:|
 | `apple` | 304 |
 | `nvidia` | 5 |
-| `rocm` | 10 |
+| `rocm` | 12 |
 | `x86` | 74 |
 
 ## Apple GPU kernel families × dtype matrix

src/compiler/codegen/Tessera_ROCM_Backend/runtime/hip/tessera_rocm_gemm.cpp

@@ -100,6 +100,75 @@ extern "C" __global__ void %NAME%(
 }
 )HIPSRC";
 
+// Rung-1 variant with K-UNROLLING: process KU 16-wide K-panels per outer step.
+// Loads KU panels of A/B (16*KU contiguous k → better memory-level parallelism /
+// coalescing) then issues KU*MT*NT WMMAs. The KU WMMAs per accumulator form a
+// short dependency chain, but the MT*NT independent chains give the scheduler
+// ILP to hide WMMA latency across a bigger loop body — vs. the extra a/b temp
+// registers (KU×) that cost occupancy on this VGPR-bound APU. Which wins is
+// EMPIRICAL (STRIX Stage H+; measured via tessera_rocm_wmma_gemm_f16_bench_ku).
+// A tail loop handles K % (16*KU). KU=1 is exactly the production kernel.
+const char* kKernelTemplateKU = R"HIPSRC(
+typedef %TYPE% wtype16 __attribute__((ext_vector_type(16)));
+typedef float  float8  __attribute__((ext_vector_type(8)));
+extern "C" __global__ void %NAME%(
+    const %TYPE%* A, const %TYPE%* B, float* D, int M, int N, int K) {
+  const int MT = %MT%, NT = %NT%, KU = %KU%;
+  int l = threadIdx.x, lane = l & 15;
+  int baseRow = blockIdx.y * 16 * MT;
+  int baseCol = blockIdx.x * 16 * NT;
+  float8 c[MT][NT];
+  for (int mi = 0; mi < MT; ++mi)
+    for (int ni = 0; ni < NT; ++ni) c[mi][ni] = (float8){0,0,0,0,0,0,0,0};
+  int k0 = 0;
+  for (; k0 + 16 * KU <= K; k0 += 16 * KU) {        // KU panels are all in-range
+    wtype16 a[MT][KU], b[NT][KU];
+    for (int u = 0; u < KU; ++u) {
+      int kb = k0 + u * 16;
+      for (int mi = 0; mi < MT; ++mi) {
+        int ar = baseRow + mi * 16 + lane;
+        bool inr = ar < M;
+        for (int i = 0; i < 16; ++i)
+          a[mi][u][i] = inr ? A[ar * K + kb + i] : (%TYPE%)0;
+      }
+      for (int ni = 0; ni < NT; ++ni) {
+        int bc = baseCol + ni * 16 + lane;
+        bool inc = bc < N;
+        for (int i = 0; i < 16; ++i)
+          b[ni][u][i] = inc ? B[(kb + i) * N + bc] : (%TYPE%)0;
+      }
+    }
+    for (int u = 0; u < KU; ++u)
+      for (int mi = 0; mi < MT; ++mi)
+        for (int ni = 0; ni < NT; ++ni)
+          c[mi][ni] = %WMMA%(a[mi][u], b[ni][u], c[mi][ni]);
+  }
+  for (; k0 < K; k0 += 16) {                         // tail: K % (16*KU)
+    wtype16 a[MT], b[NT];
+    for (int mi = 0; mi < MT; ++mi)
+      for (int i = 0; i < 16; ++i) {
+        int ar = baseRow + mi * 16 + lane, ak = k0 + i;
+        a[mi][i] = (ar < M && ak < K) ? A[ar * K + ak] : (%TYPE%)0;
+      }
+    for (int ni = 0; ni < NT; ++ni)
+      for (int i = 0; i < 16; ++i) {
+        int bk = k0 + i, bc = baseCol + ni * 16 + lane;
+        b[ni][i] = (bk < K && bc < N) ? B[bk * N + bc] : (%TYPE%)0;
+      }
+    for (int mi = 0; mi < MT; ++mi)
+      for (int ni = 0; ni < NT; ++ni)
+        c[mi][ni] = %WMMA%(a[mi], b[ni], c[mi][ni]);
+  }
+  for (int mi = 0; mi < MT; ++mi)
+    for (int ni = 0; ni < NT; ++ni)
+      for (int e = 0; e < 8; ++e) {
+        int r = baseRow + mi * 16 + e * 2 + (l >> 4);
+        int col = baseCol + ni * 16 + lane;
+        if (r < M && col < N) D[r * N + col] = c[mi][ni][e];
+      }
+}
+)HIPSRC";
+
 // Rung 2 — LDS-staged, multi-wave workgroup. A workgroup of WM×WN waves
 // cooperatively stages the A and B K-panels (16-wide) for its
 // (WM*MT*16)×(WN*NT*16) macro-tile into LDS once per K-step, then every wave
@@ -283,6 +352,19 @@ bool compileVariant(const char* type, const char* wmma, int mt, int nt,
   return compileSrc(src, name, outMod, outFn);
 }
 
+// Compile a K-unrolled rung-1 variant for (type, wmma, MT, NT, KU).
+bool compileVariantKU(const char* type, const char* wmma, int mt, int nt, int ku,
+                      const std::string& name, hipModule_t* outMod,
+                      hipFunction_t* outFn) {
+  std::string src = substitute(kKernelTemplateKU, "%TYPE%", type);
+  src = substitute(src, "%WMMA%", wmma);
+  src = substitute(src, "%NAME%", name);
+  src = substitute(src, "%MT%", std::to_string(mt));
+  src = substitute(src, "%NT%", std::to_string(nt));
+  src = substitute(src, "%KU%", std::to_string(ku));
+  return compileSrc(src, name, outMod, outFn);
+}
+
 // Compile a rung-2 LDS-staged variant for (type, wmma, WM, WN waves, MT, NT
 // register tiles/wave).
 bool compileVariantLDS(const char* type, const char* wmma, int wm, int wn,
@@ -524,6 +606,23 @@ int benchVariant(const char* type, const char* wmma, int M, int N, int K,
   return rc;
 }
 
+// K-unrolled rung-1 benchmark. 0/1/2/3 as above.
+int benchVariantKU(const char* type, const char* wmma, int M, int N, int K,
+                   int iters, int mt, int nt, int ku, double* avg_ms) {
+  if (M <= 0 || N <= 0 || K <= 0 || iters <= 0 || mt <= 0 || nt <= 0 || ku <= 0)
+    return 1;
+  hipModule_t mod = nullptr;
+  hipFunction_t fn = nullptr;
+  std::string name = std::string("benchku") + type + "_" + std::to_string(mt)
+                     + "x" + std::to_string(nt) + "_k" + std::to_string(ku);
+  if (!compileVariantKU(type, wmma, mt, nt, ku, name, &mod, &fn)) return 2;
+  unsigned gx, gy;
+  gridFor(M, N, mt, nt, &gx, &gy);
+  int rc = timedKernelLaunches(fn, gx, gy, 32, M, N, K, iters, avg_ms);
+  if (mod) hipModuleUnload(mod);
+  return rc;
+}
+
 // rung-2 (LDS-staged, WM×WN waves) benchmark. 0/1/2/3 as above.
 int benchVariantLDS(const char* type, const char* wmma, int M, int N, int K,
                     int iters, int wm, int wn, int mt, int nt, double* avg_ms) {
@@ -598,6 +697,24 @@ int runGemmPipe(const char* type, const char* wmma, const void* A,
   return rc;
 }
 
+// Run a K-unrolled GEMM end-to-end (correctness path for the KU reference rung).
+int runGemmKU(const char* type, const char* wmma, const void* A, const void* B,
+              void* D, int M, int N, int K, int mt, int nt, int ku,
+              size_t elemBytes) {
+  // Reject nonpositive tiling before gridFor (mt/nt==0 → div-by-zero) or the
+  // kernel (ku==0 → K-loop increments by 0 → GPU hang) — mirrors benchVariantKU.
+  if (M <= 0 || N <= 0 || K <= 0 || mt <= 0 || nt <= 0 || ku <= 0) return 1;
+  hipModule_t mod = nullptr;
+  hipFunction_t fn = nullptr;
+  std::string name = std::string("runku") + type;
+  if (!compileVariantKU(type, wmma, mt, nt, ku, name, &mod, &fn)) return 2;
+  unsigned gx, gy;
+  gridFor(M, N, mt, nt, &gx, &gy);
+  int rc = runDeviceGemm(fn, gx, gy, 32, A, B, D, M, N, K, elemBytes);
+  if (mod) hipModuleUnload(mod);
+  return rc;
+}
+
 // END-TO-END CPU-timed benchmark: times the full launch()-equivalent path per
 // call — memory setup + transfer + launch + sync (+ teardown) — over real host
 // buffers, for ``zerocopy`` 0 (hipMalloc + H2D/D2H copy) vs 1 (hipHostRegister
@@ -670,6 +787,15 @@ extern "C" int tessera_rocm_wmma_gemm_f16_bench(int M, int N, int K, int iters,
                       M, N, K, iters, mt, nt, avg_ms);
 }
 
+// Device-timed benchmark for the K-unrolled rung-1 kernel (f16): MT×NT register
+// tiles/wave, KU 16-wide K-panels per step. avg_ms <- mean per-launch ms.
+extern "C" int tessera_rocm_wmma_gemm_f16_bench_ku(int M, int N, int K, int iters,
+                                                   int mt, int nt, int ku,
+                                                   double* avg_ms) {
+  return benchVariantKU("__fp16", "__builtin_amdgcn_wmma_f32_16x16x16_f16_w32",
+                        M, N, K, iters, mt, nt, ku, avg_ms);
+}
+
 // Device-timed benchmark for the rung-2 LDS-staged kernel (f16): WM×WN waves per
 // workgroup, MT×NT register tiles per wave. avg_ms <- mean per-launch ms.
 extern "C" int tessera_rocm_wmma_gemm_f16_bench_lds(int M, int N, int K,
@@ -708,6 +834,14 @@ extern "C" int tessera_rocm_wmma_gemm_f16_pipe(const void* A, const void* B,
                      A, B, D, M, N, K, wm, wn, mt, nt, sizeof(unsigned short));
 }
 
+// Real-data K-unrolled GEMM (correctness path for the KU reference rung).
+extern "C" int tessera_rocm_wmma_gemm_f16_ku(const void* A, const void* B,
+                                             void* D, int M, int N, int K,
+                                             int mt, int nt, int ku) {
+  return runGemmKU("__fp16", "__builtin_amdgcn_wmma_f32_16x16x16_f16_w32",
+                   A, B, D, M, N, K, mt, nt, ku, sizeof(unsigned short));
+}
+
 // End-to-end CPU-timed benchmark of the full launch path. zerocopy: 0 = copy
 // (hipMalloc + H2D/D2H), 1 = zero-copy (hipHostRegister device-mapped). avg_ms
 // <- mean per-call ms. The APU win (if any) shows here, not in the kernel-only

tests/unit/test_rocm_gemm_ku_reference.py

@@ -0,0 +1,84 @@
+"""K-unrolled WMMA GEMM reference rung (tessera_rocm_wmma_gemm_f16_ku).
+
+The K-unroll lever — process KU 16-wide K-panels per step — was measured on
+gfx1151 and REGRESSES (the extra a/b load buffers blow the VGPR budget on this
+occupancy-bound APU; see STRIX_HALO_EXECUTION_PLAN Stage H). It is kept as a
+correctness-verified REFERENCE rung (like the LDS / pipelined rungs), not the
+production path — production stays rung-1 register blocking (size-adaptive 2x4 /
+3x4). This test locks the reference kernel's CORRECTNESS across shapes (incl.
+ragged K where the tail loop runs) and KU factors. Skip-clean w/o the .so / GPU.
+"""
+
+from __future__ import annotations
+
+import ctypes
+
+import numpy as np
+import pytest
+
+
+def _ku_or_skip():
+    from tessera import runtime as rt
+    lib = rt._load_rocm_gemm_runtime()
+    if lib is None:
+        pytest.skip("libtessera_rocm_gemm.so not loadable")
+    if not rt._rocm_wmma_runtime_available():
+        pytest.skip("no usable AMD GPU")
+    fn = getattr(lib, "tessera_rocm_wmma_gemm_f16_ku", None)
+    if fn is None:
+        pytest.skip("libtessera_rocm_gemm.so lacks the _ku entry (rebuild)")
+    fn.argtypes = [ctypes.c_void_p] * 3 + [ctypes.c_int] * 6
+    fn.restype = ctypes.c_int
+    return fn
+
+
+_RNG = np.random.default_rng(3)
+
+
+@pytest.mark.parametrize("M,N,K,mt,nt,ku", [
+    (32, 32, 32, 2, 4, 2),
+    (64, 48, 80, 3, 4, 2),
+    (48, 64, 64, 4, 3, 2),
+    (33, 17, 49, 2, 4, 2),      # ragged K (tail loop) + ragged M/N
+    (64, 64, 64, 2, 4, 4),      # KU=4
+    (80, 80, 96, 3, 4, 4),
+    (16, 16, 17, 2, 4, 2),      # K=17: one main panel + a 1-wide tail
+])
+def test_ku_reference_matches_numpy(M, N, K, mt, nt, ku):
+    fn = _ku_or_skip()
+    a = (_RNG.standard_normal((M, K)) * 0.5).astype(np.float16)
+    b = (_RNG.standard_normal((K, N)) * 0.5).astype(np.float16)
+    d = np.zeros((M, N), np.float32)
+    rc = fn(a.ctypes.data_as(ctypes.c_void_p),
+            b.ctypes.data_as(ctypes.c_void_p),
+            d.ctypes.data_as(ctypes.c_void_p), M, N, K, mt, nt, ku)
+    assert rc == 0, f"kernel rc={rc}"
+    ref = a.astype(np.float32) @ b.astype(np.float32)
+    np.testing.assert_allclose(d, ref, rtol=5e-3, atol=5e-3)
+
+
+@pytest.mark.parametrize("mt,nt,ku", [(0, 4, 2), (2, 0, 2), (2, 4, 0),
+                                       (-1, 4, 2), (2, 4, -1)])
+def test_ku_rejects_nonpositive_tiling(mt, nt, ku):
+    # mt/nt==0 would div-by-zero in gridFor; ku==0 makes the K-loop increment by
+    # 0 (GPU hang). The entry must reject (rc != 0) before launch.
+    fn = _ku_or_skip()
+    a = np.zeros((16, 16), np.float16)
+    d = np.zeros((16, 16), np.float32)
+    rc = fn(a.ctypes.data_as(ctypes.c_void_p), a.ctypes.data_as(ctypes.c_void_p),
+            d.ctypes.data_as(ctypes.c_void_p), 16, 16, 16, mt, nt, ku)
+    assert rc != 0
+
+
+def test_ku1_equals_production_semantics():
+    # KU=1 is exactly the production register-blocked kernel — must be correct.
+    fn = _ku_or_skip()
+    a = (_RNG.standard_normal((48, 64)) * 0.5).astype(np.float16)
+    b = (_RNG.standard_normal((64, 32)) * 0.5).astype(np.float16)
+    d = np.zeros((48, 32), np.float32)
+    rc = fn(a.ctypes.data_as(ctypes.c_void_p),
+            b.ctypes.data_as(ctypes.c_void_p),
+            d.ctypes.data_as(ctypes.c_void_p), 48, 32, 64, 2, 4, 1)
+    assert rc == 0
+    np.testing.assert_allclose(d, a.astype(np.float32) @ b.astype(np.float32),
+                               rtol=5e-3, atol=5e-3)