[ET-VK] Rename gemm_double_buf reference shader to coopmat_mm_ref

The test/custom_ops reference port of NVIDIA's shmem_double_buf4.comp is
the comparison baseline for the coopmat shaders, so name it after what it
is (the coopmat_mm reference) rather than how it buffers. Shader, yaml,
op (etvk.coopmat_mm_ref), impl file, bench labels, and comment references
updated; validated on device (fp16 GEMM bench: 6/6 correctness PASS,
numbers unchanged).

Authored with Claude.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

Author: xuyanwen2012

backends/vulkan/runtime/graph/ops/glsl/linear_dq8ca_qw_coopmat.glsl

@@ -27,7 +27,7 @@
  * unsigned int4 nibbles in dotPacked4x8) cancels out and is not needed.
  *
  * Loop structure follows the NVIDIA double-buffered GEMM reference
- * (shmem_double_buf4.comp "store-first" variant; see gemm_double_buf.glsl in
+ * (shmem_double_buf4.comp "store-first" variant; see coopmat_mm_ref.glsl in
  * test/custom_ops): prologue register prefetch, then per chunk
  * barrier -> prefetch next chunk -> int8 MMA on the current LDS slice ->
  * store temp into the other slice. One barrier per chunk; the prefetch is

backends/vulkan/runtime/graph/ops/glsl/linear_dq8ca_qw_coopmat.yaml

@@ -11,7 +11,7 @@
 #   WEIGHT_NBITS=8 -> linear_dq8ca_q8csw_coopmat (INT8 per-channel symmetric)
 # Requires the VK_COMPONENT_TYPE_SINT8_KHR cooperative matrix property to be
 # enumerated on the device.
-# Loop structure follows the double-buffered reference (gemm_double_buf) at
+# Loop structure follows the double-buffered reference (coopmat_mm_ref) at
 # a 128x64 tile with K-step 32, 4 subgroups x 64 threads. The reference's
 # subgroup-32 layout is NOT used — the Xclipse PAL compiler crashes in
 # vkCreateComputePipelines when int8 WMMA is compiled at forced subgroup

backends/vulkan/runtime/graph/ops/glsl/linear_qw_coopmat.glsl

@@ -22,7 +22,7 @@
  * Bsh, keeping the K-loop a clean fp16 MMA.
  *
  * Loop structure follows the NVIDIA double-buffered GEMM reference
- * (shmem_double_buf4.comp, "store-first" variant; see gemm_double_buf.glsl
+ * (shmem_double_buf4.comp, "store-first" variant; see coopmat_mm_ref.glsl
  * in test/custom_ops — measured 1.5x faster than the previous
  * single-buffered skeleton at fp16 on Xclipse 970):
  *   - PROLOGUE: prefetch tile 0 from global memory into temp registers, then
@@ -110,7 +110,7 @@ ${layout_declare_spec_const(C, "int", "num_groups_arg", "0")}
 // the first store per subgroup lands correctly; standalone repro cm_acc2).
 ${layout_declare_spec_const(C, "int", "out_N_arg", "0")}
 
-// --- Tile geometry (from yaml; defaults match gemm_double_buf) ---
+// --- Tile geometry (from yaml; defaults match coopmat_mm_ref) ---
 const uint MMA_M = ${MMA_M};
 const uint MMA_N = ${MMA_N};
 const uint MMA_K = ${MMA_K};

backends/vulkan/runtime/graph/ops/glsl/linear_qw_coopmat.yaml

@@ -12,7 +12,7 @@
 # Forces buffer storage for activation/output (coopMatLoad/Store on buffers);
 # INT weight storage can be texture2d or buffer (matches the tiled path).
 # DTYPE = half only; fp32 activations are not supported.
-# Geometry follows the double-buffered reference (gemm_double_buf): 128x128
+# Geometry follows the double-buffered reference (coopmat_mm_ref): 128x128
 # tile, K-step 16, 8 subgroups x 32 threads (subgroup size 32 forced).
 
 linear_qw_coopmat:

backends/vulkan/runtime/graph/ops/impl/QuantizedLinear.cpp

@@ -54,7 +54,7 @@ void resize_linear_qw_node(
 }
 
 // Per-shader coopmat tile geometry (must match each shader's yaml). The
-// shaders restructured to the double-buffered reference (gemm_double_buf)
+// shaders restructured to the double-buffered reference (coopmat_mm_ref)
 // use larger tiles and K-step 16; the rest keep the GemmCoopmat.h 64x64x32
 // geometry. All use 256-thread workgroups.
 //   linear_q4gsw_coopmat       128x128x16, 8 subgroups x 32 (forced)

…est/custom_ops/glsl/gemm_double_buf.glsl …test/custom_ops/glsl/coopmat_mm_ref.glslbackends/vulkan/test/custom_ops/glsl/gemm_double_buf.glsl renamed to backends/vulkan/test/custom_ops/glsl/coopmat_mm_ref.glsl backends/vulkan/test/custom_ops/glsl/gemm_double_buf.glsl renamed to backends/vulkan/test/custom_ops/glsl/coopmat_mm_ref.glsl

@@ -38,7 +38,7 @@
 
 layout(std430) buffer;
 
-// Bindings — match add_gemm_double_buf_node: output(0), mat1(1), mat2(2).
+// Bindings — match add_coopmat_mm_ref_node: output(0), mat1(1), mat2(2).
 layout(set = 0, binding = 0) buffer restrict writeonly t_outputBuffer {
     float16_t t_output[]; // fp16 D [M, N]
 };

…est/custom_ops/glsl/gemm_double_buf.yaml …test/custom_ops/glsl/coopmat_mm_ref.yamlbackends/vulkan/test/custom_ops/glsl/gemm_double_buf.yaml renamed to backends/vulkan/test/custom_ops/glsl/coopmat_mm_ref.yaml backends/vulkan/test/custom_ops/glsl/gemm_double_buf.yaml renamed to backends/vulkan/test/custom_ops/glsl/coopmat_mm_ref.yaml

@@ -8,7 +8,7 @@
 # the ET shader system. Geometry matches the reference's standalone harness:
 # one workgroup = one 128x128 output tile, K-step 16, 8 subgroups x 32 threads.
 
-gemm_double_buf:
+coopmat_mm_ref:
   parameter_names_with_default_values:
     MMA_M: 16
     MMA_N: 16
@@ -20,4 +20,4 @@ gemm_double_buf:
     SG_GRID_Y: 2
     SUBGROUP_SIZE: 32
   shader_variants:
-    - NAME: gemm_double_buf_half
+    - NAME: coopmat_mm_ref_half

…n/test/custom_ops/impl/GemmDoubleBuf.cpp …an/test/custom_ops/impl/CoopmatMmRef.cppbackends/vulkan/test/custom_ops/impl/GemmDoubleBuf.cpp renamed to backends/vulkan/test/custom_ops/impl/CoopmatMmRef.cpp backends/vulkan/test/custom_ops/impl/GemmDoubleBuf.cpp renamed to backends/vulkan/test/custom_ops/impl/CoopmatMmRef.cpp

@@ -14,26 +14,26 @@
 namespace vkcompute {
 
 // Dispatch for the ported NVIDIA double-buffered coopmat GEMM reference
-// (gemm_double_buf.glsl): D[M,N] = A[M,K] x B[K,N], all fp16 buffers,
+// (coopmat_mm_ref.glsl): D[M,N] = A[M,K] x B[K,N], all fp16 buffers,
 // row-major. One workgroup per 128x128 output tile, 256 threads, subgroup
 // size forced to 32 by the shader's REQUIRED_SUBGROUP_SIZE annotation.
 
-constexpr uint32_t kDbTileM = 128;
-constexpr uint32_t kDbTileN = 128;
-constexpr uint32_t kDbTileK = 16;
-constexpr uint32_t kDbInvocations = 256; // 8 subgroups x 32
+constexpr uint32_t kRefTileM = 128;
+constexpr uint32_t kRefTileN = 128;
+constexpr uint32_t kRefTileK = 16;
+constexpr uint32_t kRefInvocations = 256; // 8 subgroups x 32
 
-static vkapi::ShaderInfo pick_gemm_double_buf_shader(
+static vkapi::ShaderInfo pick_coopmat_mm_ref_shader(
     ComputeGraph* graph,
     const std::vector<ArgGroup>& args,
     const std::vector<ValueRef>& resize_args) {
   (void)graph;
   (void)args;
   (void)resize_args;
-  return VK_KERNEL_FROM_STR("gemm_double_buf_half");
+  return VK_KERNEL_FROM_STR("coopmat_mm_ref_half");
 }
 
-static utils::uvec3 pick_gemm_double_buf_global_wg_size(
+static utils::uvec3 pick_coopmat_mm_ref_global_wg_size(
     ComputeGraph* graph,
     const vkapi::ShaderInfo& shader,
     const std::vector<ArgGroup>& args,
@@ -45,15 +45,15 @@ static utils::uvec3 pick_gemm_double_buf_global_wg_size(
   const uint32_t M = out_sizes.at(out_sizes.size() - 2);
   const uint32_t N = out_sizes.at(out_sizes.size() - 1);
   // Same group-count cancellation trick as GemmCoopmat.cpp: the framework
-  // divides by the local size, so multiplying tiles_n by kDbInvocations
+  // divides by the local size, so multiplying tiles_n by kRefInvocations
   // yields exactly tiles_n x tiles_m workgroups.
   return {
-      utils::div_up(N, kDbTileN) * kDbInvocations,
-      utils::div_up(M, kDbTileM),
+      utils::div_up(N, kRefTileN) * kRefInvocations,
+      utils::div_up(M, kRefTileM),
       1};
 }
 
-static utils::uvec3 pick_gemm_double_buf_local_wg_size(
+static utils::uvec3 pick_coopmat_mm_ref_local_wg_size(
     ComputeGraph* graph,
     const vkapi::ShaderInfo& shader,
     const utils::uvec3& global_workgroup_size,
@@ -64,10 +64,10 @@ static utils::uvec3 pick_gemm_double_buf_local_wg_size(
   (void)global_workgroup_size;
   (void)args;
   (void)resize_args;
-  return {kDbInvocations, 1, 1};
+  return {kRefInvocations, 1, 1};
 }
 
-void gemm_double_buf(ComputeGraph& graph, const std::vector<ValueRef>& args) {
+void coopmat_mm_ref(ComputeGraph& graph, const std::vector<ValueRef>& args) {
   int idx = 0;
   const ValueRef mat1 = args.at(idx++);
   const ValueRef mat2 = args.at(idx++);
@@ -82,15 +82,15 @@ void gemm_double_buf(ComputeGraph& graph, const std::vector<ValueRef>& args) {
   const int32_t N = graph.size_at<int32_t>(-1, out);
   const int32_t K = graph.size_at<int32_t>(-1, mat1);
   // No partial-tile or K-tail handling in the reference shader.
-  VK_CHECK_COND(M % static_cast<int32_t>(kDbTileM) == 0);
-  VK_CHECK_COND(N % static_cast<int32_t>(kDbTileN) == 0);
-  VK_CHECK_COND(K % static_cast<int32_t>(kDbTileK) == 0);
+  VK_CHECK_COND(M % static_cast<int32_t>(kRefTileM) == 0);
+  VK_CHECK_COND(N % static_cast<int32_t>(kRefTileN) == 0);
+  VK_CHECK_COND(K % static_cast<int32_t>(kRefTileK) == 0);
 
   graph.execute_nodes().emplace_back(new DynamicDispatchNode(
       graph,
-      pick_gemm_double_buf_shader,
-      pick_gemm_double_buf_global_wg_size,
-      pick_gemm_double_buf_local_wg_size,
+      pick_coopmat_mm_ref_shader,
+      pick_coopmat_mm_ref_global_wg_size,
+      pick_coopmat_mm_ref_local_wg_size,
       // Inputs and Outputs
       {{out, vkapi::kWrite}, {{mat1, mat2}, vkapi::kRead}},
       // Shader params buffers — none; all geometry is spec constants
@@ -106,7 +106,7 @@ void gemm_double_buf(ComputeGraph& graph, const std::vector<ValueRef>& args) {
 }
 
 REGISTER_OPERATORS {
-  VK_REGISTER_OP(etvk.gemm_double_buf, gemm_double_buf);
+  VK_REGISTER_OP(etvk.coopmat_mm_ref, coopmat_mm_ref);
 }
 
 } // namespace vkcompute

backends/vulkan/test/custom_ops/test_fp16_gemm_bench.cpp

@@ -8,7 +8,7 @@
 //   tiled    matmul_vec, Texture3D (production default baseline)
 //   coopmat  matmul_coopmat (coopmat_mm.glsl), Buffer — our shader, forced
 //            past the desktop-only gate via test_etvk.test_mm "coopmat"
-//   dbuf4    gemm_double_buf (NVIDIA shmem_double_buf4 reference port),
+//   coopmat_ref  coopmat_mm_ref (NVIDIA shmem_double_buf4 reference port),
 //            Buffer — double-buffered shared memory, subgroup size 32
 //
 // Apples-to-apples: same shapes, same runtime-mat2 row-major [K,N] fp16
@@ -30,9 +30,9 @@ struct GemmConfig {
   int64_t N;
 };
 
-// Impl rows benchmarked per shape. The dbuf4 tile is 128x128 (vs 64x64 for
+// Impl rows benchmarked per shape. The coopmat_ref tile is 128x128 (vs 64x64 for
 // coopmat), so correctness shapes must align to 128.
-static const std::vector<std::string> kImpls = {"tiled", "coopmat", "dbuf4"};
+static const std::vector<std::string> kImpls = {"tiled", "coopmat", "coopmat_ref"};
 
 static TestCase make_case(const GemmConfig& cfg, const std::string& impl) {
   const vkapi::ScalarType dt = vkapi::kHalf;
@@ -53,8 +53,8 @@ static TestCase make_case(const GemmConfig& cfg, const std::string& impl) {
   ValueSpec output({cfg.M, cfg.N}, dt, storage, utils::kWidthPacked,
                    DataGenType::ZEROS);
 
-  if (impl == "dbuf4") {
-    tc.set_operator_name("etvk.gemm_double_buf");
+  if (impl == "coopmat_ref") {
+    tc.set_operator_name("etvk.coopmat_mm_ref");
     tc.add_input_spec(mat1);
     tc.add_input_spec(mat2);
   } else {
@@ -68,7 +68,7 @@ static TestCase make_case(const GemmConfig& cfg, const std::string& impl) {
   }
   tc.add_output_spec(output);
 
-  // tiled accumulates in fp16 (error grows with K); coopmat/dbuf4 accumulate
+  // tiled accumulates in fp16 (error grows with K); coopmat/coopmat_ref accumulate
   // in fp32, bounded by fp16 input/output rounding only.
   if (impl == "tiled") {
     tc.set_abs_tolerance(1.0f);
@@ -128,8 +128,8 @@ std::vector<TestCase> generate_cases() {
       }
     }
   }
-  // Correctness: aligned to the dbuf4 128x128 tile (and coopmat's 64/32);
-  // the second shape dispatches a 2x2 workgroup grid for dbuf4.
+  // Correctness: aligned to the coopmat_ref 128x128 tile (and coopmat's 64/32);
+  // the second shape dispatches a 2x2 workgroup grid for coopmat_ref.
   static const std::vector<GemmConfig> kCorrectnessShapes = {
       {128, 64, 128}, {256, 128, 256}};
   for (const auto& cfg : kCorrectnessShapes) {
@@ -172,7 +172,7 @@ int main() {
     for (const auto& st : r.get_shader_timings()) {
       if (st.shader_name.find("matmul") != std::string::npos ||
           st.shader_name.find("coopmat") != std::string::npos ||
-          st.shader_name.find("double_buf") != std::string::npos) {
+          st.shader_name.find("coopmat_mm_ref") != std::string::npos) {
         name = st.shader_name;
       }
     }
@@ -183,7 +183,7 @@ int main() {
             << ") ==========\n";
   std::cout << std::left << std::setw(15) << "shape(K,N)" << std::right
             << std::setw(10) << "tiled" << std::setw(10) << "coopmat"
-            << std::setw(10) << "dbuf4" << std::setw(12) << "dbuf4/coop"
+            << std::setw(12) << "coopmat_ref" << std::setw(10) << "ref/coop"
             << "  kernels\n";
   size_t idx = 0;
   for (const auto& kn : kShapes) {
@@ -201,7 +201,7 @@ int main() {
                   std::to_string(kn.second) + ")")
               << std::right << std::fixed << std::setprecision(1)
               << std::setw(10) << tiled << std::setw(10) << coop
-              << std::setw(10) << dbuf << std::setw(11)
+              << std::setw(12) << dbuf << std::setw(9)
               << std::setprecision(2) << (coop > 0 ? dbuf / coop : 0.0f)
               << "x  " << c_kernel << " | " << d_kernel << "\n";
   }