perf(cuda): release MMA stages earlier and flatten native weight tiles

humming/config/config.py

@@ -80,6 +80,7 @@ class LayerConfig(BaseHummingConfig):
         "is_token_input_scale_2",
         "is_tensor_input_scale_2",
         "use_raw_weight",
+        "use_flat_weight",
         "use_native_dequant",
         "use_block_scaled_mma",
         "mma_scale_group_size",
@@ -132,7 +133,7 @@ def is_umma_supported(self):
     @property
     def use_raw_weight(self):
         if self.a_dtype.num_bits == self.b_dtype.num_bits:
-            return True
+            return self.is_umma_supported or current_device.is_ppu
         has_native_mixed_activation = self.a_dtype in (
             dtypes.float8e4m3,
             dtypes.float8e5m2,
@@ -145,6 +146,11 @@ def use_raw_weight(self):
         )
         return has_native_mixed_activation and has_native_mixed_weight and self.is_umma_supported
 
+    @property
+    def use_flat_weight(self):
+        has_same_width = self.a_dtype.num_bits == self.b_dtype.num_bits
+        return has_same_width and not self.is_umma_supported and not current_device.is_ppu
+
     @property
     def use_block_scaled_mma(self):
         if self.mxmma_supported:
@@ -168,7 +174,7 @@ def mma_scale_group_size(self):
     @property
     def use_native_dequant(self):
         assert self.sm_version is not None
-        if self.use_raw_weight or self.sm_version // 10 == 9:
+        if self.use_raw_weight or self.use_flat_weight or self.sm_version // 10 == 9:
             return False
 
         from humming.jit.runtime import KernelRuntime
@@ -379,6 +385,7 @@ def __post_init__(self):
             can_use_packed_k = (
                 self.sm_version // 10 == 9
                 and not self.use_raw_weight
+                and not self.use_flat_weight
                 and has_packed_k_operands
                 and has_packed_k_input_scale
             )
@@ -389,7 +396,9 @@ def __post_init__(self):
             self.use_packed_k_layout = can_use_packed_k and prefer_packed_k
         elif self.use_packed_k_layout:
             assert self.sm_version // 10 == 9, "use_packed_k_layout requires SM90"
-            assert not self.use_raw_weight, "raw weights do not use packed-K layout"
+            assert not (self.use_raw_weight or self.use_flat_weight), (
+                "raw/flat weights do not use packed-K layout"
+            )
             assert self.a_dtype.num_bits == 8, "use_packed_k_layout requires 8-bit activation"
             assert self.b_dtype.num_bits % 2 == 0, "use_packed_k_layout requires even-bit weight"
 

humming/config/mma.py

@@ -174,10 +174,11 @@ def get_register_budget_error(layer_config, tuning_config, use_f16_accum=False,
         return None  # UMMA accumulators use TMEM and have separate resource checks.
 
     if mma_type == MmaType.WGMMA:
+        use_wgmma_ss = layer_config.a_dtype.num_bits == layer_config.b_dtype.num_bits
         # ptxas needs 26 (SS) or 32 (RS) registers beyond one N=64 accumulator tuple.
         # setmaxnreg cannot relax this per-instruction compile limit.
         instruction_accumulators = warp_m // (4 if use_f16_accum else 2)
-        instruction_overhead = 26 if layer_config.use_raw_weight else 32
+        instruction_overhead = 26 if use_wgmma_ss else 32
         instruction_registers = instruction_accumulators + instruction_overhead
         if layer_config.use_packed_k_layout and tuning_config.wgmma_split_issue_wait:
             # Split packed-K keeps all RS K slabs live across prefetch. The

humming/csrc/launcher/launcher.cpp

@@ -336,6 +336,7 @@ std::tuple<int64_t, std::string> register_kernel(const std::string &cubin_path)
       reader.getBool("USE_PACKED_K_LAYOUT"),
       reader.getBool("USE_UMMA_SS"),
       reader.getBool("USE_RAW_WEIGHT"),
+      reader.getBool("USE_FLAT_WEIGHT"),
       reader.getBool("USE_BLOCK_SCALED_MMA")};
 
   std::unique_lock lock(g_kernel_mutex);

humming/csrc/launcher/tensor.h

@@ -87,7 +87,7 @@ inline void check_tensor_b(Tensor &tensor, KernelData &kernel_data, int64_t dev)
 
   std::vector<int64_t> expected_shape = {};
   if (kernel_data.gemm_type_id != 0) expected_shape.push_back(kernel_data.num_experts);
-  if (kernel_data.use_raw_weight) {
+  if (kernel_data.use_raw_weight || kernel_data.use_flat_weight) {
     if (get_dtype_num_bits(kernel_data.a_dtype_id) == 8 && num_bits < 8)
       problem_shape_k = CEIL_DIV(problem_shape_k, 128) * 128;
     expected_shape.push_back(problem_shape_n);
@@ -390,6 +390,24 @@ inline CUtensorMap make_tma_desc_b(Tensor &tensor, KernelData &kernel_data) {
   uint32_t block_shape_n = kernel_data.block_shape_n;
   uint32_t block_shape_k = kernel_data.block_shape_k;
 
+  if (kernel_data.use_flat_weight) {
+    CUtensorMap descriptor{};
+    uint64_t warp_n = kernel_data.mma_type == MmaType::WGMMA ? kernel_data.warp_shape_n : 16;
+    uint64_t shape_n = kernel_data.problem_shape_n;
+    uint64_t num_experts = std::max(1u, kernel_data.num_experts);
+    uint64_t row_bytes = uint64_t(tensor.size(-1)) * 4;
+    uint64_t dimensions[] = {64, 64 / warp_n, shape_n / 64, warp_n / 16, row_bytes / 16 * num_experts};
+    uint64_t strides[] = {warp_n * 16, 1024, 256, shape_n * 16};
+    uint32_t box[] = {64, uint32_t(std::min(64u, block_shape_n) / warp_n),
+        (block_shape_n + 63) / 64, uint32_t(warp_n / 16), block_shape_k * num_bits / 128};
+    uint32_t element_strides[] = {1, 1, 1, 1, 1};
+    CUresult status = cuTensorMapEncodeTiled(&descriptor, CU_TENSOR_MAP_DATA_TYPE_INT32, 5,
+        tensor.data_ptr(), dimensions, strides, box, element_strides, CU_TENSOR_MAP_INTERLEAVE_NONE,
+        CU_TENSOR_MAP_SWIZZLE_NONE, CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
+    ASSERT_CHECK(status == CUDA_SUCCESS, "TMA Encode Failed for flat weights: ", int(status));
+    return descriptor;
+  }
+
   if (kernel_data.use_raw_weight) {
     // FP4/FP6 in an f8f6f4 instruction uses padded 16-element blocks in SMEM.
     bool expand_low_bit = num_bits < 8 && get_dtype_num_bits(kernel_data.a_dtype_id) == 8;
@@ -408,25 +426,6 @@ inline CUtensorMap make_tma_desc_b(Tensor &tensor, KernelData &kernel_data) {
       ASSERT_CHECK(status == CUDA_SUCCESS, "TMA Encode Failed for expanded low-bit weights");
       return descriptor;
     }
-    if (kernel_data.mma_type == MmaType::WGMMA) {
-      CUtensorMap descriptor{};
-      uint64_t w = kernel_data.warp_shape_n;
-      uint64_t row_bytes = uint64_t(tensor.size(-1)) * 4;
-      // Fragment planes span all N warps in the CTA, leaving the fifth
-      // dimension available for K slabs even with multiple N warpgroups.
-      uint64_t dimensions[] = {swizzle_bytes / 4, 16,
-          uint64_t(kernel_data.problem_shape_n) * std::max(1u, kernel_data.num_experts) / w,
-          w / 16, row_bytes / swizzle_bytes};
-      uint64_t strides[] = {row_bytes, w * row_bytes, 16 * row_bytes, swizzle_bytes};
-      uint32_t box[] = {swizzle_bytes / 4, 16, uint32_t(block_shape_n / w), uint32_t(w / 16),
-          block_shape_k * num_bits / 8 / swizzle_bytes};
-      uint32_t element_strides[] = {1, 1, 1, 1, 1};
-      CUresult status = cuTensorMapEncodeTiled(&descriptor, CU_TENSOR_MAP_DATA_TYPE_INT32, 5,
-          tensor.data_ptr(), dimensions, strides, box, element_strides, CU_TENSOR_MAP_INTERLEAVE_NONE,
-          get_swizzle_enum(swizzle_bytes), CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
-      ASSERT_CHECK(status == CUDA_SUCCESS, "TMA Encode Failed for WGMMA N permutation: ", int(status));
-      return descriptor;
-    }
     auto rows = torch_view_shape(tensor, {-1, tensor.size(-1)});
     return make_tma_desc(rows, {swizzle_bytes / 4, std::min(256u, kernel_data.use_umma_ss ? std::max(128u, block_shape_n) : block_shape_n)}, swizzle_bytes, "b");
   }

humming/csrc/launcher/utils.h

@@ -179,6 +179,7 @@ struct KernelData {
   bool use_packed_k_layout;
   bool use_umma_ss;
   bool use_raw_weight;
+  bool use_flat_weight;
   bool use_block_scaled_mma;
 };
 

humming/include/humming/kernel/humming_ws.cuh

@@ -189,6 +189,9 @@ __global__ __launch_bounds__(TuningConfig::kNumThreads, TuningConfig::kNumCtasPe
           if (warp_iter_id == Ctx::kWarpIters - 1 && slice_iter + 1 < num_slice_iters) {
             consumer.wait_stage((stage_id + 1) % kNumStages);
           }
+          if constexpr (!Ctx::kUseWgmma) {
+            if (warp_iter_id == Ctx::kWarpIters - 1) consumer.arrive(stage_id);
+          }
           const bool use_late_scale_prefetch = Ctx::kUseWgmmaLateScalePrefetch && warp_iter_id == Ctx::kWarpIters - 1;
           const bool has_next_fragment = warp_iter_id + 1 < Ctx::kWarpIters || slice_iter + 1 < num_slice_iters;
           if (has_next_fragment) {
@@ -208,7 +211,9 @@ __global__ __launch_bounds__(TuningConfig::kNumThreads, TuningConfig::kNumCtasPe
                 ((warp_iter_id + 1) % Ctx::kWarpIters) % 2,
                 (warp_iter_id + 1) % Ctx::kWarpIters);
           }
-          if (warp_iter_id == Ctx::kWarpIters - 1) consumer.arrive(stage_id);
+          if constexpr (Ctx::kUseWgmma) {
+            if (warp_iter_id == Ctx::kWarpIters - 1) consumer.arrive(stage_id);
+          }
         }
       };
 

humming/include/humming/memory/g2s_loader/loader_b.cuh

@@ -54,7 +54,8 @@ public:
 
   template <bool kShouldAdvance = true>
   CUDA_INLINE void load(int4 *smem_ptr, void *mbar_ptr) {
-    if constexpr (Ctx::kUseRawWeight && !kUseTma) load_raw(smem_ptr);
+    if constexpr (Ctx::kUseFlatWeight && !kUseTma) load_flat(smem_ptr);
+    else if constexpr (Ctx::kUseRawWeight && !kUseTma) load_raw(smem_ptr);
     else if constexpr (kUseAiu) load_aiu(smem_ptr);
     else if constexpr (kUseTma) load_tma(smem_ptr, mbar_ptr);
     else load_legacy(smem_ptr);
@@ -76,22 +77,21 @@ public:
   CUDA_INLINE
   void load_tma(int4 *smem_ptr, void *mbar_ptr) {
     if (ctx.load_thread_id() == 0) {
-      if constexpr (Ctx::kUseRawWeight) {
+      if constexpr (Ctx::kUseFlatWeight) {
+        uint32_t expert_k = col_offset / ProblemShape::N * ProblemShape::K + row_offset;
+        if (cluster_rank == 0)
+          tma_load_5d<kMultiCastSizeB>(tensor_map_ptr, smem_ptr, mbar_ptr,
+              0, col_offset % 64 / 16, col_offset % ProblemShape::N / 64, 0, expert_k * ElementQB::kBits / 128);
+      } else if constexpr (Ctx::kUseRawWeight) {
         constexpr uint32_t kCoordinateBits = SharedStorage::kExpandUmmaWeight ? 1 : 32 / ElementQB::kBits;
         uint32_t start_k = kCoordinateBits == 1 ? row_offset / 128 * 128 : row_offset;
-        if constexpr (Ctx::kUseWgmmaTmaKPack) {
-          if (cluster_rank == 0)
-            tma_load_5d<kMultiCastSizeB>(tensor_map_ptr, smem_ptr, mbar_ptr,
-                                       0, 0, col_offset / Ctx::WarpShape::N, 0, start_k / kSwizzleK);
-        } else {
+        PRAGMA_UNROLL
+        for (uint32_t k = 0; k < SharedStorage::kWeightStageK; k += kSwizzleK) {
           PRAGMA_UNROLL
-          for (uint32_t k = 0; k < SharedStorage::kWeightStageK; k += kSwizzleK) {
-            PRAGMA_UNROLL
-            for (uint32_t n = 0; n < SharedStorage::kWeightStageN; n += 256) {
-              uint32_t offset = (k / kSwizzleK * SharedStorage::kWeightStageN + n) * kSwizzleBytes / sizeof(int4);
-              if (cluster_rank == 0) tma_load_2d<kMultiCastSizeB, kEvictWeightsFirst, Ctx::kUseUmmaCooperativeTma ? 2 : 1>(
-                  tensor_map_ptr, smem_ptr + offset, mbar_ptr, (start_k + k) / kCoordinateBits, col_offset + n);
-            }
+          for (uint32_t n = 0; n < SharedStorage::kWeightStageN; n += 256) {
+            uint32_t offset = (k / kSwizzleK * SharedStorage::kWeightStageN + n) * kSwizzleBytes / sizeof(int4);
+            if (cluster_rank == 0) tma_load_2d<kMultiCastSizeB, kEvictWeightsFirst, Ctx::kUseUmmaCooperativeTma ? 2 : 1>(
+                tensor_map_ptr, smem_ptr + offset, mbar_ptr, (start_k + k) / kCoordinateBits, col_offset + n);
           }
         }
       } else if constexpr (kMultiCastSizeB == 1) {
@@ -106,12 +106,12 @@ public:
   void prefetch_tma() {
     if constexpr (kUseTma && kMultiCastSizeB == 1) {
       if (ctx.load_thread_id() == 0) {
-        if constexpr (!Ctx::kUseRawWeight) tma_prefetch_3d(tensor_map_ptr, 0, col_offset, row_offset);
+        if constexpr (!Ctx::kUseRawWeight && !Ctx::kUseFlatWeight) tma_prefetch_3d(tensor_map_ptr, 0, col_offset, row_offset);
       }
     }
   }
 
-  CUDA_INLINE static uint32_t permute_raw_ss_row(uint32_t row) {
+  CUDA_INLINE static uint32_t permute_wgmma_ss_row(uint32_t row) {
     constexpr uint32_t kWarpN = Ctx::WarpShape::N;
     constexpr uint32_t kGroupN = 4 * kWarpN;
     uint32_t row_in_group = row % kGroupN;
@@ -127,36 +127,34 @@ public:
     constexpr uint32_t kGlobalRowInt4s = ProblemShape::K * ElementQB::kBits / 128;
     constexpr uint32_t kSwizzleInt4s = kSwizzleBytes / sizeof(int4);
     uint32_t smem_base = cast_smem_ptr_to_uint(smem_ptr) / 128;
-    constexpr uint32_t kRowsPerIteration = kNumLoadThreads / kRowInt4s;
-    constexpr bool kPowerOfTwoRows = kRowsPerIteration != 0 && (kRowsPerIteration & (kRowsPerIteration - 1)) == 0;
-    constexpr bool kPowerOfTwoWarpN = (Ctx::WarpShape::N & (Ctx::WarpShape::N - 1)) == 0;
-    constexpr bool kSplitRowPermutation =
-        Ctx::kUseWgmmaCpAsyncNPermute && kNumLoadThreads % kRowInt4s == 0 && kPowerOfTwoRows && kPowerOfTwoWarpN;
-    uint32_t permuted_thread_row = 0;
-    if constexpr (kSplitRowPermutation) {
-      // The thread and unrolled iteration own disjoint row bits. Permute
-      // thread bits once; the iteration contribution folds to constants.
-      uint32_t thread_row = (ctx.load_thread_id() / kRowInt4s) % kRowsPerIteration;
-      permuted_thread_row = permute_raw_ss_row(thread_row);
-    }
     PRAGMA_UNROLL
     for (uint32_t i = 0; i < CEIL_DIV(BlockShape::N * kRowInt4s, kNumLoadThreads); ++i) {
       uint32_t index = i * kNumLoadThreads + ctx.load_thread_id();
       uint32_t row = index / kRowInt4s;
       uint32_t col = index % kRowInt4s;
-      uint32_t smem_row = row;
-      if constexpr (Ctx::kUseWgmmaCpAsyncNPermute) {
-        if constexpr (kSplitRowPermutation) {
-          smem_row = permuted_thread_row + permute_raw_ss_row(i * kRowsPerIteration);
-        } else smem_row = permute_raw_ss_row(row);
-      }
-      uint32_t offset = (col / kSwizzleInt4s * BlockShape::N + smem_row) * kSwizzleInt4s + col % kSwizzleInt4s;
+      uint32_t offset = (col / kSwizzleInt4s * BlockShape::N + row) * kSwizzleInt4s + col % kSwizzleInt4s;
       uint32_t swizzled_offset = offset ^ ((smem_base + offset / 8) % kSwizzleInt4s);
       legacy_load_pred<kUseCpAsync>(gmem_ptr + row * kGlobalRowInt4s + col,
                                     smem_ptr + swizzled_offset, row < BlockShape::N);
     }
   }
 
+  CUDA_INLINE
+  void load_flat(int4 *smem_ptr) {
+    constexpr uint32_t kChunkRows = MIN(64u, BlockShape::N);
+    constexpr uint32_t kTileInt4s = kChunkRows * BlockShape::K * ElementQB::kBits / 128;
+    PRAGMA_UNROLL
+    for (uint32_t i = 0; i < CEIL_DIV(kNumInt4s, kNumLoadThreads); ++i) {
+      uint32_t index = i * kNumLoadThreads + ctx.load_thread_id();
+      uint32_t row = index / kTileInt4s * 64 + index % kChunkRows;
+      uint32_t col = index % kTileInt4s / kChunkRows;
+      if constexpr (Ctx::kUseWgmmaCpAsyncNPermute) row = permute_wgmma_ss_row(row);
+      uint32_t offset = col * BlockShape::N + row;
+      uint32_t gmem_offset = col * ProblemShape::N + index / kTileInt4s * 64 + index % kChunkRows;
+      legacy_load_pred<kUseCpAsync>(gmem_ptr + gmem_offset, smem_ptr + offset, index < kNumInt4s);
+    }
+  }
+
   CUDA_INLINE
   void load_raw_aiu(int4 *smem_ptr) {
     const uint32_t warp_id = ctx.load_thread_id() / 32;
@@ -179,9 +177,9 @@ public:
 
   CUDA_INLINE
   void advance() {
-    if constexpr (Ctx::kUseRawWeight) {
+    if constexpr (Ctx::kUseRawWeight || Ctx::kUseFlatWeight) {
       row_offset += BlockShape::K;
-      if constexpr (!kUseTma) gmem_ptr += BlockShape::K * ElementQB::kBits / 128;
+      if constexpr (!kUseTma) gmem_ptr += BlockShape::K * ElementQB::kBits / 128 * (Ctx::kUseFlatWeight ? ProblemShape::N : 1);
       return;
     }
     row_offset += BlockShape::K / kPackSizeK;
@@ -190,12 +188,18 @@ public:
 
   CUDA_INLINE
   void seek(uint32_t expert_id, uint32_t n_block_id, uint32_t k_block_id) {
-    if constexpr (Ctx::kUseRawWeight) {
+    if constexpr (Ctx::kUseRawWeight || Ctx::kUseFlatWeight) {
       row_offset = k_block_id * BlockShape::K;
       col_offset = expert_id * ProblemShape::N + n_block_id * BlockShape::N;
       if constexpr (!kUseTma) {
-        uint64_t offset = uint64_t(col_offset) * ProblemShape::K + row_offset;
-        gmem_ptr = gmem_ptr_raw + offset * ElementQB::kBits / 128;
+        if constexpr (Ctx::kUseFlatWeight) {
+          uint64_t expert_offset = uint64_t(expert_id) * ProblemShape::N * ProblemShape::K;
+          uint64_t offset = expert_offset + uint64_t(row_offset) * ProblemShape::N;
+          gmem_ptr = gmem_ptr_raw + offset * ElementQB::kBits / 128 + n_block_id * BlockShape::N;
+        } else {
+          uint64_t offset = uint64_t(col_offset) * ProblemShape::K + row_offset;
+          gmem_ptr = gmem_ptr_raw + offset * ElementQB::kBits / 128;
+        }
       }
       return;
     }

humming/include/humming/memory/s2r_loader/loader_b.cuh

@@ -62,7 +62,7 @@ public:
   }
 
   CUDA_INLINE
-  void load_raw(const int4 *smem_ptr, uint32_t *regs_ptr, uint32_t iter_id) {
+  void load_native(const int4 *smem_ptr, uint32_t *regs_ptr, uint32_t iter_id) {
     if constexpr (USE_PPU && Ctx::kUseCpAsync && !Ctx::kUseTmaB)
       return load_raw_aiu(smem_ptr, regs_ptr, iter_id);
     constexpr uint32_t kSwizzleInt4s = MIN(8u, BlockShape::K * ElementQB::kBits / 128);
@@ -76,6 +76,7 @@ public:
       uint32_t col = (ctx.k_warp_offset() * ElementQB::kBits / 256 + iter_id) * 2 + (lane_id / 8) % 2;
       uint32_t offset = (col / kSwizzleInt4s * BlockShape::N + row) * kSwizzleInt4s + col % kSwizzleInt4s;
       uint32_t swizzled_offset = offset ^ ((smem_base + offset / 8) % kSwizzleInt4s);
+      if constexpr (Ctx::kUseFlatWeight) swizzled_offset = col * BlockShape::N + row;
       ld_shared<kRowsPerLoad / 4>(smem_ptr + swizzled_offset, reinterpret_cast<int4 *>(regs_ptr + n * kRowsPerLoad / 4));
     }
   }
@@ -94,7 +95,7 @@ public:
 
   CUDA_INLINE
   void load(const int4 *smem_ptr, uint32_t *regs_ptr, uint32_t iter_id) {
-    if constexpr (Ctx::kUseRawWeight) return load_raw(smem_ptr, regs_ptr, iter_id);
+    if constexpr (Ctx::kUseRawWeight || Ctx::kUseFlatWeight) return load_native(smem_ptr, regs_ptr, iter_id);
     if constexpr (Ctx::kUsePackedKLayout) return load_packed_k(smem_ptr, regs_ptr, iter_id);
     uint32_t warp_id = ctx.warp_id();
     uint32_t n_warp_id = ctx.n_warp_id();

humming/include/humming/memory/s2r_pipeline.cuh

@@ -85,7 +85,7 @@ public:
     uint32_t buffer_id = iter_id % 2;
     auto &smem = ctx.smem;
 
-    if constexpr (!(kUseWgmma && Ctx::kUseRawWeight))
+    if constexpr (!Ctx::kUseWgmmaSs)
       loader_b.load(smem.stages[stage_id].b, mma.regs_qb_as_ptr(buffer_id), iter_id);
     if constexpr (!kUseWgmma && !Ctx::kUseUmma)
       loader_a.load(smem.stages[stage_id].a, mma.regs_a_as_ptr(buffer_id), iter_id, stage_id);

humming/include/humming/mma/mxmma.cuh

@@ -70,8 +70,7 @@ public:
 
   CUDA_INLINE
   void transform_b(uint32_t buffer_id, uint32_t iter_id) {
-    if constexpr (Ctx::kUseRawWeight) {
-      // Raw operands are loaded directly into regs_b, including mixed FP4 formats.
+    if constexpr (Ctx::kUseRawWeight || Ctx::kUseFlatWeight) {
       return;
     } else if constexpr (kNativeMixed) {
       PRAGMA_UNROLL
@@ -180,7 +179,7 @@ public:
 
   template <class T = uint32_t>
   CUDA_INLINE T *regs_qb_as_ptr(uint32_t buffer_id) {
-    if constexpr (Ctx::kUseRawWeight) {
+    if constexpr (Ctx::kUseRawWeight || Ctx::kUseFlatWeight) {
       return reinterpret_cast<T *>(regs_b[buffer_id]);
     } else {
       return reinterpret_cast<T *>(regs_qb[buffer_id]);