review(mk-gemm v2): PR #305 리뷰 13건 반영 -- 스윕 기반 ksr 규칙·장치 SM 수·정확 역수·복사 언롤·정적 에필로그, 벤치 기본 레인, 트레이스 도구 메인 스트림; 원장 30차 §4~5

MEASUREMENTS.md

@@ -2020,9 +2020,70 @@ partial 은 m × n × ksr ≤ 32 × 6528 × 4. 벤치 `--gemm2 both`(v1 열 옆
 기대(측정 전): 메인 99발 6.4 → ~4.9 ms(고정비 ~3 µs, 균형 꼬리 제거), aux 쌍 노출 28 →
 ~5 µs/층(down 이 MoE 꼬리 안에서 끝남) → 스텝 −2.5 ms, 트레이스 mk_gemm 합 14.1 → ~6.5~7.
 
-### 4. 측정 (채워질 자리)
-
-_srv2 `~/mkab/g2.bench.out` / `g2.overlap.out` — 플릿이 lever-chain4 로 점유돼 GPU 창 대기._
+### 4. 측정 1차 — 단독 벤치 (2026-09-05 00:48, srv2, chain4 와 chain5 사이의 3분 창, `g2.bench.out`, iters 10)
+
+수치 게이트: exact grid v1 3.9e-8 / v2 ksr 0·1·2·3·5·8 전부 ≤ 3.9e-8, over-ulp 0, replay 안정 —
+**PASS**. v2 출력은 비분할 형상(k=512, 4096², k=2048, k=3072, m=16 4096²)에서 v1 과 **비트
+동일**, 분할 형상은 슬라이스 합 순서 차이 1e-10~5e-6.
+
+단독(콜드 L2)/연속 2발(x2) µs, m=8 이 C=1:
+
+| 형상 | stock | v1 | v1 x2 | v2(규칙) | v2 x2 | 규칙 ksr/유닛 | 스윕 최선 (ksr: 단독/x2) |
+|---|---|---|---|---|---|---|---|
+| [6416×4096] in_proj | 133 | 84.0 | 80.9 | 93.2 | 86.0 | 4 / 204 | 8: 86.0/79.4 |
+| [4096×4096] MLA o | 100 | 61.4 | 58.4 | 60.4 | 54.8 | 6 / 192 | 2: 55.3/55.3 |
+| [4096×2048] KDA o | 49 | 34.8 | 31.7 | 41.0 | 34.8 | 4 / 128 | 2: 34.8/32.8 |
+| [2048×4096] MLA qkv | 49 | 36.9 | 32.8 | 39.9 | 34.8 | 8 / 128 | 6: 35.8/29.7 |
+| [6144×4096] dense gu | 142 | 78.7 | 74.8 | 80.9 | 73.6 | 4 / 192 | 2: 74.8/70.7 |
+| [4096×3072] dense dn | 73 | 50.2 | 46.6 | 50.2 | 44.0 | 6 / 192 | 2: 45.1/44.0 |
+| [1024×4096] 공유 gu | 34 | 26.6 | 22.5 | **24.6** | **21.5** | 8 / 64 | 8 |
+| [4096×512] 공유 dn | 20.5 | 18.4 | 15.4 | **16.4** | **14.8** | 1 / 32 | 1: 16.4/13.9 |
+| m=16 [4096×4096] | 97 | 63.5 | 60.4 | 82.9 | 67.9 | 6 / 192 | 2: 63.5/62.5 |
+| m=32 [2048×4096] | 54 | 41.9 | 37.9 | 53.2 | 49.6 | 8 / 128 | 6: 45.0/40.9 |
+| m=32 [1024×4096] | 35 | 30.7 | 28.6 | 33.8 | 31.2 | 8 / 64 | 8: 33.0/30.7 |
+
+읽기: (1) 공유 전문가 두 형상은 단독으로도 v1 보다 빠르다(−8~11%) — 서빙에서 보는 것은
+이 값이 아니라 MoE 옆의 노출(§2)이고 그건 노출 프로브가 잰다. (2) 큰 형상은 **ksr 규칙이
+틀렸다**: "2 웨이브 + 슬라이스 ≥ 4" 가 고른 값(4·6·8)보다 스윕의 최선은 유닛 수가 상주 슬롯
+96 에 딱 맞거나(6144 × 2 = 96, 2048 × 6 = 96) 그 아래로 긴 슬라이스(4096: ksr 2 = 64 유닛)
+인 쪽이고, 51 타일(in_proj)만 잘게(ksr 8 = 408) 썰어야 꼬리가 줄었다(ksr 2 = 102 유닛은
+6개짜리 두 번째 웨이브가 +30 µs). 최선을 쓰면 v2 는 m=8 전 형상에서 v1 과 같거나 빠르다
+(6416 86/79 vs 84/81, 4096² 55 vs 61, 6144 75/71 vs 79/75, 3072 45/44 vs 50/47).
+(3) **m ≥ 16 은 v2 가 느리다**(m=16 4096² 83 vs 63.5, m=32 2048 53 vs 42): 블록마다 k-블록당
+x 를 다시 읽고 양자화하는 비용이 m 에 비례하고(warp 당 4행), 그 L2 적재 지연이 mma 뒤에
+드러난다 — v1 은 발사당 한 번이었다. C=1(m=8)에는 없는 비용이지만 C≥2 브래킷 전에 고쳐야
+한다(후보: m ≤ 16 특화 인스턴스 + x 2단 선적재, 또는 k-블록별 first-arriver 양자화 공유).
+
+### 5. 측정 2차 — ksr 3 스윕과 규칙 확정 (2026-09-05 01:11, srv2, chain5 뒤 창, `g2.ksr3.out`)
+
+| 형상 (m=8) | ksr 2 | ksr 3 | ksr 4 | ksr 8 | v1 x2 |
+|---|---|---|---|---|---|
+| [6416×4096] 51타일 | 112.6/98.8 | 87.0/82.4 | 93.2/86.6 | **86.0/79.4** | 80.4 |
+| [4096×4096] 32 | 56.3/55.3 | **55.3/49.7** | 66.6/61.9 | 61.4/56.3 | 56.3 |
+| [4096×2048] 32 | 34.8/32.8 | **35.8/29.7** | 40.9/34.8 | 37.9/34.8 | 32.8 |
+| [2048×4096] 16 | 43.0/41.0 | 34.8/33.8 | **34.8/32.8** | 41.0/34.8 | 33.2 |
+| [6144×4096] 48 | **76.0/70.7** | 85.0/82.6 | 79.0/73.2 | 81.9/75.8 | 72.8 |
+| [4096×3072] 32 | 45.1/44.5 | **44.1/39.9** | 55.3/48.1 | 51.2/45.6 | 46.1 |
+| m=16 [4096×4096] | 62.5/62.0 | **56.3/51.7** | 70.7/65.1 | 68.6/62.0 | 59.4 |
+| m=32 [2048×4096] | 61.4/58.4 | **47.1/45.1** | 47.1/45.6 | 53.3/48.6 | 37.9 |
+
+(단독/x2 µs.) **유닛 수 = 상주 슬롯 96 딱 한 웨이브**(48×2, 32×3, 16×6)가 k 와 무관하게
+최선이고, 51 타일은 짧은 두 번째 웨이브(102)가 최악이라 잘게(8) 썰어야 하며, 8 타일은
+슬라이스 ≥ 4 k-블록 안에서 최대(8 = 64 유닛). 이 규칙을 `mk_choose_ksr2` 에 넣었다(SM 수는
+장치에서). 규칙 적용 시 v2 x2 vs v1 x2: in_proj 79.4/80.4, 4096² 49.7/56.3, k2048 29.7/32.8,
+n2048 32.8/33.2, 6144 70.7/72.8, k3072 39.9/46.1, 1024 21.5/22.5, k512 14.8/15.4, m=16 4096²
+51.7/59.4 — **m ≤ 16 전 형상에서 v2 ≥ v1**. m=32 는 45.1 vs 37.9 로 여전히 뒤진다(§4 (3)).
+
+노출 프로브 1차는 스톡 flashinfer 래퍼의 `run()` 에 `out=` 이 없어 실패(벤치 러너는
+메가커널 파일만 마운트하고, `out=` 은 배포 b12x 오버레이의 확장) — 반환값을 쓰도록 고쳐
+재대기. 유닛 스탬프도 다음 창.
+
+**리뷰(PR #305 머지 뒤, 13건)에서 나온 커널 쪽 사실**(SASS 판독): k-블록당 warp 명령 ~670
+중 LUT 쌍 재구성(프래그먼트마다 16쌍) 216, 스테이지 복사 루프 미언롤 80, 행당 양자화 70
+(`1.0f/sc` 가 IEEE 나눗셈 슬로패스 호출), 에필로그의 런타임 바운드 때문에 acc 가 로컬
+메모리(64 B 스택). 후속 PR 에서 복사 언롤·`__frcp_rn`(2의 거듭제곱이라 정확)·정적 인덱스를
+넣었고, LUT 쌍 공유(k 순열로 프래그먼트 두 개를 한 그룹에)는 v1 과의 비트 동일을 깨므로
+스탬프로 발행 한계가 확인되면 다음 라운드.
 
 ## ★★★28차 — 드래프터 W4 는 서빙된 적이 없었다(컴파일 캐시), MK-MLA 서빙 사망의 원인은 스크래치 재할당, 그리고 드래프터 W4 판정 (2026-09-04)
 

RUNBOOK_KERNEL_CAMPAIGN2.md

@@ -595,9 +595,11 @@ deep_gemm 은 독립 블록이라 같은 GEMM 을 MoE 꼬리 안에서 끝냈다
 두 레인 × ksr 5종).
 
 **오프라인 게이트(부팅 전, srv2 GPU 창)**: `run_megakernel_bench.sh --segments exact,gemm
---gemm2 both --ksr2-sweep 1,2,4,6,8` 전 행 PASS + v2 열이 v1 열보다 형상마다 빠름(특히
-[6144×4096]·[4096×512]·[1024×4096]) + `mk_gemm_moe_overlap_probe.py` 의 v2 노출이 v1 의
-28 µs/층보다 작음. 상한: 스텝 −2.5 ms(메인 −1.5, aux 노출 −1.0), 트레이스 mk 합 14.1 → ~7.
+--gemm2 both`(기본 `--gemm2 env` 는 서빙 레인만 판정한다) 전 행 PASS + v2 열이 v1 열보다
+형상마다 같거나 빠름 + `mk_gemm_moe_overlap_probe.py` 의 v2 노출이 v1 의 28 µs/층보다 작음.
+상태(09-05 01:11, 원장 30차 §4~5): exact PASS(두 레인 × ksr 5종), 규칙 확정 뒤 m ≤ 16 전
+형상 v2 ≥ v1(x2), m = 32 는 뒤짐(C=1 전용 팔) — 노출 프로브·스탬프 대기. 상한: 스텝 −2.5 ms
+(메인 −1.5, aux 노출 −1.0), 트레이스 mk 합 14.1 → ~7.
 
 **부팅 게이트**: base(기본값) → cand(`VLLM_GLM53_MK_GEMM2=1`) 브래킷, step/s(acc 정규화)
 + 프리필 동반 + 품질 9/9 + 한국어 0/16 + pos-1 ±2 pct. 통과하면 프로필 기본 1 로 올리고

build/dsv4/glm53_megakernel.cu

@@ -1017,6 +1017,14 @@ __global__ void mk_gemm_kernel(const MKGemmCtx c) {
 // expanded tiles. ksr > 1 slices go to an fp32 partial and the last slice
 // to arrive folds the tile in fixed slice order (the persistent kernel's
 // leftover fold, applied to every tile): deterministic, no zero pass.
+//
+// Shared state: the partials and the per-tile arrival counters below are
+// one set for the whole device, like the persistent lane's A tiles and
+// unit counter -- so, like every MK GEMM launch, a v2 launch must never
+// overlap another MK GEMM launch on a different stream (two launches
+// folding the same tile index would sum each other's slices). Serving
+// keeps that by stream order: the side-stream pair is joined before the
+// next main-stream GEMM.
 // ===========================================================================
 // v2 raw-record ring depth (VLLM_GLM53_MK_NBUF2; 2..4 keep two blocks per SM)
 #ifndef MK_NBUF2_DEF
@@ -1112,7 +1120,12 @@ mk_gemm2_kernel(const MKGemm2Ctx c) {
     const uint8_t* ssrc = (const uint8_t*)c.ws4 +
         ((size_t)nt * kblk + kb) * (SMEM_W_ROWS * 8);
     uint8_t* d = sraw + buf * W4_RAW_BYTES;
-    for (int t = (int)threadIdx.x; t < SMEM_W_ROWS * 4; t += MK_THREADS) {
+    // two chunks per thread, unrolled: the swizzled destinations are
+    // per-thread constants (a runtime loop recomputed them every k-block)
+    static_assert((SMEM_W_ROWS * 4) % MK_THREADS == 0, "chunks per thread");
+#pragma unroll
+    for (int u = 0; u < (SMEM_W_ROWS * 4) / MK_THREADS; ++u) {
+      const int t = (int)threadIdx.x + u * MK_THREADS;
       const int r = t >> 2, ch = t & 3;
       mk_cp_async16(d + r * W4_RAW_PITCH + ((ch ^ ((r >> 1) & 3)) << 4),
                     nsrc + (size_t)t * 16);
@@ -1153,7 +1166,10 @@ mk_gemm2_kernel(const MKGemm2Ctx c) {
       for (int off = 16; off; off >>= 1)
         mx = fmaxf(mx, __shfl_xor_sync(0xffffffffu, mx, off));
       const float sc = mk_pow2_scale(mx);
-      const float rsc = 1.0f / sc;  // exact: sc is a power of two
+      // sc is a power of two, so rcp.rn is exact -- the same rsc as the
+      // shared prologue's 1.0f / sc, without the IEEE divide's slow-path
+      // call that the SASS showed on every row of every k-block
+      const float rsc = __frcp_rn(sc);
       uint32_t pack = 0;
 #pragma unroll
       for (int q = 0; q < 4; ++q)
@@ -1289,51 +1305,32 @@ mk_gemm2_kernel(const MKGemm2Ctx c) {
   }
   MK2_TS(2);
 
-  // ---- epilogue
-  if (ksr == 1) {  // whole tile: bf16 out, real rows / cols only
+  // ---- epilogue: one walk over the fragment's real rows / cols, two stores
+  auto store_tile = [&](auto&& put) {  // put(row, col, value)
 #pragma unroll
-    for (int i = 0; i < mtiles; ++i) {
+    for (int i = 0; i < 2; ++i) {  // static index: a runtime bound put acc in local memory
+      if (i >= mtiles) break;
 #pragma unroll
       for (int j = 0; j < 2; ++j) {
         const int r0 = i * 16 + g, r1 = r0 + 8;
         const int cb = nt * 128 + warp * 16 + j * 8 + (lane & 3) * 2;
-        if (r0 < c.m) {
-          if (cb < c.n_orig)
-            c.out[(size_t)r0 * c.n_orig + cb] = __float2bfloat16(acc[i][j][0]);
-          if (cb + 1 < c.n_orig)
-            c.out[(size_t)r0 * c.n_orig + cb + 1] = __float2bfloat16(acc[i][j][1]);
-        }
-        if (r1 < c.m) {
-          if (cb < c.n_orig)
-            c.out[(size_t)r1 * c.n_orig + cb] = __float2bfloat16(acc[i][j][2]);
-          if (cb + 1 < c.n_orig)
-            c.out[(size_t)r1 * c.n_orig + cb + 1] = __float2bfloat16(acc[i][j][3]);
-        }
+        if (r0 < c.m) { put(r0, cb, acc[i][j][0]); put(r0, cb + 1, acc[i][j][1]); }
+        if (r1 < c.m) { put(r1, cb, acc[i][j][2]); put(r1, cb + 1, acc[i][j][3]); }
       }
     }
+  };
+  if (ksr == 1) {  // whole tile: bf16 out
+    store_tile([&](int r, int col, float v) {
+      if (col < c.n_orig) c.out[(size_t)r * c.n_orig + col] = __float2bfloat16(v);
+    });
     MK2_TS(3);
     return;
   }
   // k-slice: assign (never accumulate) this slice's partial, count the
   // arrival, and let the last slice fold the tile in slice order.
   {
     float* pb = g_mk2_partial + (size_t)sp * c.m * c.n;
-#pragma unroll
-    for (int i = 0; i < mtiles; ++i) {
-#pragma unroll
-      for (int j = 0; j < 2; ++j) {
-        const int r0 = i * 16 + g, r1 = r0 + 8;
-        const int cb = nt * 128 + warp * 16 + j * 8 + (lane & 3) * 2;
-        if (r0 < c.m) {
-          pb[(size_t)r0 * c.n + cb] = acc[i][j][0];
-          pb[(size_t)r0 * c.n + cb + 1] = acc[i][j][1];
-        }
-        if (r1 < c.m) {
-          pb[(size_t)r1 * c.n + cb] = acc[i][j][2];
-          pb[(size_t)r1 * c.n + cb + 1] = acc[i][j][3];
-        }
-      }
-    }
+    store_tile([&](int r, int col, float v) { pb[(size_t)r * c.n + col] = v; });
   }
   __syncthreads();
   __threadfence();  // release: the slice is visible device-wide first
@@ -2664,6 +2661,7 @@ bool g_attrs_set = false;
 // resident blocks per SM the device reports for mk_gemm2_kernel (2 by
 // construction of GEMM2_SMEM; the v2 unit rule sizes its grid from it)
 int g_gemm2_bps = 0;
+int g_mk_sms = 0;  // multiprocessors, from the device (48 on GB10)
 
 void set_kernel_attrs() {
   if (g_attrs_set) return;
@@ -2680,6 +2678,8 @@ void set_kernel_attrs() {
   MK_CHECK_CUDA(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
       &g_gemm2_bps, mk_gemm2_kernel, MK_THREADS, GEMM2_SMEM));
   if (g_gemm2_bps < 1) g_gemm2_bps = 1;
+  MK_CHECK_CUDA(cudaDeviceGetAttribute(
+      &g_mk_sms, cudaDevAttrMultiProcessorCount, 0));
   MK_CHECK_CUDA(cudaFuncSetAttribute(
       mk_mla_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, MLA_SMEM));
   g_attrs_set = true;
@@ -2797,12 +2797,18 @@ bool mk_gemm2_on() {
   }
   return g_probe_gemm2 == 1;
 }
-// k-slices per tile for one v2 launch. Enough units for two waves of the
-// resident slots (blocks/SM x 48 SMs), so the scheduler has something to
-// balance the tail with, but no slice shorter than 4 k-blocks (a unit's
-// ring fill and quant are paid once per unit). Then the contract clamps:
-// ksr <= kblk (every slice non-empty), ksr <= MK2_KSR_MAX, and the fp32
-// partial must fit (m x n x ksr floats).
+// k-slices per tile for one v2 launch, from the 30차 sweeps (srv2, ksr 1/2/3/
+// 4/6/8 on every production shape, single and back-to-back): what wins is
+// ONE exact wave of the resident slots (blocks/SM x SMs = 96 on GB10) --
+// 48 tiles x 2, 32 x 3, 16 x 6 measured best on every k -- or, when the
+// tile count does not divide the slots, the longest slices that still fit
+// one wave (8 tiles x 8 = 64 units); a small SECOND wave is the worst case
+// (51 tiles x 2 = 102 units: +30 us for the six stragglers), so a tile
+// count above half the slots takes the finest slices instead (51 x 8 = 408,
+// 86 us; x 3 = 153, 87; x 1 = 51, 97). No slice shorter than 4 k-blocks (a
+// unit's ring fill and first quant are paid per unit; k = 512 stays whole).
+// Then the contract clamps: ksr <= kblk (every slice non-empty), ksr <=
+// MK2_KSR_MAX, and the fp32 partial must fit (m x n x ksr floats).
 int mk_choose_ksr2(int m, int n, int k) {
   const int nblk = n / SMEM_W_ROWS, kblk = k / KSTEP;
   if (g_probe_ksr2 < 0) {
@@ -2813,10 +2819,17 @@ int mk_choose_ksr2(int m, int n, int k) {
   if (g_probe_ksr2 > 0) {
     ksr = g_probe_ksr2;
   } else {
-    const int slots = (g_gemm2_bps > 0 ? g_gemm2_bps : 2) * 48;
-    ksr = (2 * slots + nblk - 1) / nblk;
-    const int kmax = kblk / 4;
-    if (ksr > kmax) ksr = kmax;
+    const int slots = (g_gemm2_bps > 0 ? g_gemm2_bps : 2) *
+                      (g_mk_sms > 0 ? g_mk_sms : 48);
+    const int kmax = kblk / 4 > 1 ? kblk / 4 : 1;
+    if (slots % nblk == 0 && slots / nblk <= kmax) {
+      ksr = slots / nblk;               // one exact wave
+    } else if (nblk * 2 > slots) {
+      ksr = kmax;                       // would leave a short second wave: slice fine
+    } else {
+      ksr = slots / nblk;               // under one wave, the longest slices
+      if (ksr > kmax) ksr = kmax;
+    }
   }
   if (ksr < 1) ksr = 1;
   if (ksr > kblk) ksr = kblk;
@@ -2840,18 +2853,27 @@ std::vector<int64_t> mk_probe_device() {
           (int64_t)prop.sharedMemPerBlockOptin};
 }
 
-// Phase timestamps of the last gemm launch, [MK_GRID_CAP][8] ns, then
-// cleared. Empty unless built with -DMK_PHASE_TS=1.
-std::vector<int64_t> mk_read_ts() {
 #ifdef MK_PHASE_TS
-  std::vector<unsigned long long> h(MK_GRID_CAP * 8);
+// Read a device stamp array into a host vector and clear it (every probe
+// reader below): synchronize first -- the stamps are written by the launch
+// the caller just issued.
+template <size_t N>
+std::vector<int64_t> mk_read_and_clear(unsigned long long (&sym)[N]) {
+  std::vector<unsigned long long> h(N);
   MK_CHECK_CUDA(cudaDeviceSynchronize());
-  MK_CHECK_CUDA(cudaMemcpyFromSymbol(h.data(), g_mk_ts,
-                                     sizeof(unsigned long long) * h.size()));
+  MK_CHECK_CUDA(cudaMemcpyFromSymbol(h.data(), sym, sizeof(unsigned long long) * N));
   void* p = nullptr;
-  MK_CHECK_CUDA(cudaGetSymbolAddress(&p, g_mk_ts));
-  MK_CHECK_CUDA(cudaMemset(p, 0, sizeof(unsigned long long) * h.size()));
+  MK_CHECK_CUDA(cudaGetSymbolAddress(&p, sym));
+  MK_CHECK_CUDA(cudaMemset(p, 0, sizeof(unsigned long long) * N));
   return std::vector<int64_t>(h.begin(), h.end());
+}
+#endif
+
+// Phase timestamps of the last gemm launch, [MK_GRID_CAP][8] ns, then
+// cleared. Empty unless built with -DMK_PHASE_TS=1.
+std::vector<int64_t> mk_read_ts() {
+#ifdef MK_PHASE_TS
+  return mk_read_and_clear(g_mk_ts);
 #else
   return {};
 #endif
@@ -2861,28 +2883,15 @@ std::vector<int64_t> mk_read_ts() {
 // p2 done, barrier, p3 done, barrier, p4 done.
 std::vector<int64_t> mk_read_kda_ts() {
 #ifdef MK_PHASE_TS
-  std::vector<unsigned long long> h(MK_GRID_CAP * 16);
-  MK_CHECK_CUDA(cudaMemcpyFromSymbol(h.data(), g_mk_kda_ts,
-                                     sizeof(unsigned long long) * h.size()));
-  void* p = nullptr;
-  MK_CHECK_CUDA(cudaGetSymbolAddress(&p, g_mk_kda_ts));
-  MK_CHECK_CUDA(cudaMemset(p, 0, sizeof(unsigned long long) * h.size()));
-  return std::vector<int64_t>(h.begin(), h.end());
+  return mk_read_and_clear(g_mk_kda_ts);
 #else
   return {};
 #endif
 }
 
 std::vector<int64_t> mk_read_mhc_ts() {
 #ifdef MK_PHASE_TS
-  std::vector<unsigned long long> h(MK_MHC_GRID_CAP * 8);
-  MK_CHECK_CUDA(cudaDeviceSynchronize());
-  MK_CHECK_CUDA(cudaMemcpyFromSymbol(h.data(), g_mk_mhc_ts,
-                                     sizeof(unsigned long long) * h.size()));
-  void* p = nullptr;
-  MK_CHECK_CUDA(cudaGetSymbolAddress(&p, g_mk_mhc_ts));
-  MK_CHECK_CUDA(cudaMemset(p, 0, sizeof(unsigned long long) * h.size()));
-  return std::vector<int64_t>(h.begin(), h.end());
+  return mk_read_and_clear(g_mk_mhc_ts);
 #else
   return {};
 #endif
@@ -3119,14 +3128,7 @@ void mk_set_gemm2(int64_t on, int64_t ksr) {
 // v2 unit timestamps of the last launch, [MK2_UNITS_MAX][4] ns, then cleared.
 std::vector<int64_t> mk_read_ts2() {
 #ifdef MK_PHASE_TS
-  std::vector<unsigned long long> h(MK2_UNITS_MAX * 4);
-  MK_CHECK_CUDA(cudaDeviceSynchronize());
-  MK_CHECK_CUDA(cudaMemcpyFromSymbol(h.data(), g_mk2_ts,
-                                     sizeof(unsigned long long) * h.size()));
-  void* p = nullptr;
-  MK_CHECK_CUDA(cudaGetSymbolAddress(&p, g_mk2_ts));
-  MK_CHECK_CUDA(cudaMemset(p, 0, sizeof(unsigned long long) * h.size()));
-  return std::vector<int64_t>(h.begin(), h.end());
+  return mk_read_and_clear(g_mk2_ts);
 #else
   return {};
 #endif