diff --git a/csrc/cpu/sgl-kernels/common.h b/csrc/cpu/sgl-kernels/common.h
index de1260bb33..dc0769450e 100644
--- a/csrc/cpu/sgl-kernels/common.h
+++ b/csrc/cpu/sgl-kernels/common.h
@@ -280,26 +280,10 @@ inline int get_thread_num() {
 // balance payload across each thread
 template <typename T>
 inline void balance211(T n, T nth, T ith, T& n_start, T& n_end) {
-#if 0
-    // onednn partition pattern
-    T& n_my = n_end;
-    if (nth <= 1 || n == 0) {
-        n_start = 0;
-        n_my = n;
-    } else {
-        T n1 = div_up(n, nth);
-        T n2 = n1 - 1;
-        T T1 = n - n2 * nth;
-        n_my = ith < T1 ? n1 : n2;
-        n_start = ith <= T1 ? ith*n1 : T1 * n1 + (ith - T1) * n2;
-    }
-    n_end += n_start;
-#else
   // pytorch aten partition pattern
   T n_my = div_up(n, nth);
   n_start = ith * n_my;
   n_end = std::min(n_start + n_my, n);
-#endif
 }
 
 template <typename func_t>
diff --git a/csrc/libtorch_stable/cuda_vec_utils.cuh b/csrc/libtorch_stable/cuda_vec_utils.cuh
index ec6e60724e..c962c99ec2 100644
--- a/csrc/libtorch_stable/cuda_vec_utils.cuh
+++ b/csrc/libtorch_stable/cuda_vec_utils.cuh
@@ -320,10 +320,6 @@ __host__ __device__ __forceinline__ bool is_16byte_aligned(const void* ptr) {
   return (reinterpret_cast<uintptr_t>(ptr) & 15) == 0;
 }
 
-__host__ __device__ __forceinline__ bool is_32byte_aligned(const void* ptr) {
-  return (reinterpret_cast<uintptr_t>(ptr) & 31) == 0;
-}
-
 // ============================================================
 // Packed type conversion and arithmetic
 // ============================================================
diff --git a/csrc/libtorch_stable/moe/moe_wna16_utils.h b/csrc/libtorch_stable/moe/moe_wna16_utils.h
index 8ef03f0e60..6ce52ee5b1 100644
--- a/csrc/libtorch_stable/moe/moe_wna16_utils.h
+++ b/csrc/libtorch_stable/moe/moe_wna16_utils.h
@@ -34,10 +34,6 @@ class ScalarType<half> {
   static __host__ __device__ float2 inline num22float2(const half2 x) {
     return __half22float2(x);
   }
-
-  static __host__ __device__ half2 inline float22num2(const float2 x) {
-    return __float22half2_rn(x);
-  }
 };
 
 template <>
@@ -71,10 +67,6 @@ class ScalarType<nv_bfloat16> {
   static __host__ __device__ float2 inline num22float2(const nv_bfloat162 x) {
     return __bfloat1622float2(x);
   }
-
-  static __host__ __device__ nv_bfloat162 inline float22num2(const float2 x) {
-    return __float22bfloat162_rn(x);
-  }
 #endif
 };
 
diff --git a/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh b/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh
index 3d82e43931..abe3368cb1 100644
--- a/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh
+++ b/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh
@@ -74,18 +74,9 @@ class MatrixView_half_rw {
                                                 const int width)
       : data(data), height(height), width(width) {}
 
-  __device__ __forceinline__ half item(int row, int column) const {
-    return data[row * width + column];
-  }
   __device__ __forceinline__ half* item_ptr(int row, int column) {
     return &data[row * width + column];
   }
-  __device__ __forceinline__ void set(int row, int column, half value) {
-    data[row * width + column] = value;
-  }
-  __device__ __forceinline__ void set_half2(int row, int column, half2 value) {
-    ((half2*)data)[(row * width + column) / 2] = value;
-  }
 
   __device__ __forceinline__ void set4(int row, int column, half v0, half v1,
                                        half v2, half v3) {
@@ -113,14 +104,6 @@ class MatrixView_q4_row {
     return (data[row * width / 8 + column / 8] >> shift) & 0x0f;
   }
 
-  __device__ __forceinline__ void item2(int (&items)[2], int row,
-                                        int column) const {
-    int shift = (column & 0x07) * 4;
-    uint32_t d = data[row * width / 8 + column / 8] >> shift;
-    items[0] = d & 0x0f;
-    items[1] = (d >> 4) & 0x0f;
-  }
-
   __device__ __forceinline__ void item4(int (&items)[4], int row,
                                         int column) const {
     int shift = (column & 0x07) * 4;
@@ -148,14 +131,6 @@ class MatrixView_q2_row {
     return (data[row * width / 16 + column / 16] >> shift) & 0x03;
   }
 
-  __device__ __forceinline__ void item2(int (&items)[2], int row,
-                                        int column) const {
-    int shift = (column & 0x0f) * 2;
-    uint32_t d = data[row * width / 16 + column / 16] >> shift;
-    items[0] = d & 0x03;
-    items[1] = (d >> 2) & 0x03;
-  }
-
   __device__ __forceinline__ void item4(int (&items)[4], int row,
                                         int column) const {
     int shift = (column & 0x0f) * 2;
@@ -237,14 +212,6 @@ class MatrixView_q8_row {
     return (data[row * width / 4 + column / 4] >> shift) & 0xff;
   }
 
-  __device__ __forceinline__ void item2(int (&items)[2], int row,
-                                        int column) const {
-    int shift = (column & 0x03) * 8;
-    uint32_t d = data[row * width / 4 + column / 4] >> shift;
-    items[0] = d & 0xff;
-    items[1] = (d >> 8) & 0xff;
-  }
-
   __device__ __forceinline__ void item4(int (&items)[4], int row,
                                         int column) const {
     int shift = (column & 0x03) * 2;
diff --git a/csrc/libtorch_stable/quantization/marlin/marlin.cuh b/csrc/libtorch_stable/quantization/marlin/marlin.cuh
index bfb65e874b..af4d4730ba 100644
--- a/csrc/libtorch_stable/quantization/marlin/marlin.cuh
+++ b/csrc/libtorch_stable/quantization/marlin/marlin.cuh
@@ -67,14 +67,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr,
   }
 }
 
-__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr,
-                                         bool pred = true) {
-  if (pred) {
-    reinterpret_cast<int4*>(smem_ptr)[0] =
-        reinterpret_cast<const int4*>(glob_ptr)[0];
-  }
-}
-
 __device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr,
                                       bool pred = true) {
   if (pred) {
@@ -121,19 +113,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr,
       "r"(smem), "l"(glob_ptr), "n"(BYTES));
 }
 
-__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr,
-                                         bool pred = true) {
-  const int BYTES = 16;
-  uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
-  asm volatile(
-      "{\n"
-      "   .reg .pred p;\n"
-      "   setp.ne.b32 p, %0, 0;\n"
-      "   @p cp.async.ca.shared.global [%1], [%2], %3;\n"
-      "}\n" ::"r"((int)pred),
-      "r"(smem), "l"(glob_ptr), "n"(BYTES));
-}
-
 __device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr,
                                       bool pred = true) {
   const int BYTES = 16;
@@ -170,4 +149,4 @@ __device__ inline void cp_async_wait() {
 
 }  // namespace MARLIN_NAMESPACE_NAME
 
-#endif
\ No newline at end of file
+#endif
diff --git a/csrc/libtorch_stable/sampler.cu b/csrc/libtorch_stable/sampler.cu
index b8d64f7af1..703a748ce6 100644
--- a/csrc/libtorch_stable/sampler.cu
+++ b/csrc/libtorch_stable/sampler.cu
@@ -45,11 +45,6 @@ __global__ void apply_repetition_penalties_kernel(
   }
 }
 
-__device__ __forceinline__ auto convert_to_uint32(float x) -> uint32_t {
-  uint32_t bits = __float_as_uint(x);
-  return (bits & 0x80000000) ? bits : ~bits & 0x7fffffff;
-}
-
 template <int step>
 static inline __device__ uint32_t extractBinIdx(float x) {
   if constexpr (step == 0) {
diff --git a/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh b/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh
index 58d19de734..32925c0c3c 100644
--- a/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh
+++ b/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh
@@ -67,104 +67,6 @@ using fp8_type = __hip_fp8_e4m3_fnuz;
 using fp8x2_type = __hip_fp8x2_e4m3_fnuz;
     #endif
 
-// fp8 -> half
-template <>
-__inline__ __device__ uint16_t
-vec_conversion<uint16_t, uint8_t>(const uint8_t& a) {
-  return __hip_cvt_fp8_to_halfraw(a, fp8_type::__default_interpret).x;
-}
-
-// fp8x2 -> half2
-template <>
-__inline__ __device__ uint32_t
-vec_conversion<uint32_t, uint16_t>(const uint16_t& a) {
-  union {
-    __half2_raw h2r;
-    uint32_t ui32;
-  } tmp;
-  tmp.h2r = __hip_cvt_fp8x2_to_halfraw2(a, fp8_type::__default_interpret);
-  return tmp.ui32;
-}
-
-// fp8x4 -> half2x2
-template <>
-__inline__ __device__ uint2 vec_conversion<uint2, uint32_t>(const uint32_t& a) {
-  union {
-    uint2 u32x2;
-    uint32_t u32[2];
-  } tmp;
-  tmp.u32[0] = vec_conversion<uint32_t, uint16_t>((uint16_t)a);
-  tmp.u32[1] = vec_conversion<uint32_t, uint16_t>((uint16_t)(a >> 16U));
-  return tmp.u32x2;
-}
-
-// fp8x8 -> half2x4
-template <>
-__inline__ __device__ uint4 vec_conversion<uint4, uint2>(const uint2& a) {
-  union {
-    uint4 u64x2;
-    uint2 u64[2];
-  } tmp;
-  tmp.u64[0] = vec_conversion<uint2, uint32_t>(a.x);
-  tmp.u64[1] = vec_conversion<uint2, uint32_t>(a.y);
-  return tmp.u64x2;
-}
-
-using __nv_bfloat16 = __hip_bfloat16;
-
-// fp8 -> __nv_bfloat16
-template <>
-__inline__ __device__ __nv_bfloat16
-vec_conversion<__nv_bfloat16, uint8_t>(const uint8_t& a) {
-  fp8_type f8;
-  f8.__x = a;
-  return __float2bfloat16(static_cast<float>(f8));
-}
-
-using __nv_bfloat162 = __hip_bfloat162;
-
-// fp8x2 -> __nv_bfloat162
-template <>
-__inline__ __device__ __nv_bfloat162
-vec_conversion<__nv_bfloat162, uint16_t>(const uint16_t& a) {
-  __nv_bfloat162 res;
-  res.x = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)a);
-  res.y = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)(a >> 8U));
-  return res;
-}
-
-// fp8x4 -> bf16_4_t
-template <>
-__inline__ __device__ bf16_4_t
-vec_conversion<bf16_4_t, uint32_t>(const uint32_t& a) {
-  bf16_4_t res;
-  res.x = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)a);
-  res.y = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)(a >> 16U));
-  return res;
-}
-
-// fp8x8 -> bf16_8_t
-template <>
-__inline__ __device__ bf16_8_t vec_conversion<bf16_8_t, uint2>(const uint2& a) {
-  bf16_4_t tmp1, tmp2;
-  tmp1 = vec_conversion<bf16_4_t, uint32_t>(a.x);
-  tmp2 = vec_conversion<bf16_4_t, uint32_t>(a.y);
-  bf16_8_t res;
-  res.x = tmp1.x;
-  res.y = tmp1.y;
-  res.z = tmp2.x;
-  res.w = tmp2.y;
-  return res;
-}
-
-// fp8 -> float
-template <>
-__inline__ __device__ float vec_conversion<float, uint8_t>(const uint8_t& a) {
-  fp8_type f8;
-  f8.__x = a;
-  return static_cast<float>(f8);
-}
-
 // fp8x2 -> float2
 template <>
 __inline__ __device__ float2
@@ -193,138 +95,6 @@ vec_conversion<float4, uint32_t>(const uint32_t& a) {
   return res;
 }
 
-// fp8x8 -> float8
-template <>
-__inline__ __device__ Float8_ vec_conversion<Float8_, uint2>(const uint2& a) {
-  Float4_ tmp1, tmp2;
-  tmp1 = vec_conversion<Float4_, uint32_t>(a.x);
-  tmp2 = vec_conversion<Float4_, uint32_t>(a.y);
-  Float8_ res;
-  res.x = tmp1.x;
-  res.y = tmp1.y;
-  res.z = tmp2.x;
-  res.w = tmp2.y;
-  return res;
-}
-
-// half -> fp8
-template <>
-__inline__ __device__ uint8_t
-vec_conversion<uint8_t, uint16_t>(const uint16_t& a) {
-  __half_raw tmp;
-  tmp.x = a;
-  return __hip_cvt_halfraw_to_fp8(tmp, fp8_type::__default_saturation,
-                                  fp8_type::__default_interpret);
-}
-
-template <>
-__inline__ __device__ uint16_t
-vec_conversion<uint16_t, uint32_t>(const uint32_t& a) {
-  union {
-    uint32_t ui32;
-    __half2_raw h2r;
-  } tmp;
-  tmp.ui32 = a;
-  return __hip_cvt_halfraw2_to_fp8x2(tmp.h2r, fp8_type::__default_saturation,
-                                     fp8_type::__default_interpret);
-}
-
-// bf16 -> fp8
-template <>
-__inline__ __device__ uint8_t
-vec_conversion<uint8_t, __nv_bfloat16>(const __nv_bfloat16& a) {
-  return __hip_cvt_float_to_fp8(__bfloat162float(a),
-                                fp8_type::__default_saturation,
-                                fp8_type::__default_interpret);
-}
-
-// float -> fp8
-template <>
-__inline__ __device__ uint8_t vec_conversion<uint8_t, float>(const float& a) {
-  return __hip_cvt_float_to_fp8(a, fp8_type::__default_saturation,
-                                fp8_type::__default_interpret);
-}
-
-// float2 -> half2
-template <>
-__inline__ __device__ uint32_t
-vec_conversion<uint32_t, float2>(const float2& a) {
-  union {
-    half2 float16;
-    uint32_t uint32;
-  };
-
-  float16 = __float22half2_rn(a);
-  return uint32;
-}
-
-// Float4 -> half2x2
-template <>
-__inline__ __device__ uint2 vec_conversion<uint2, Float4_>(const Float4_& a) {
-  uint2 b;
-  float2 val;
-  val.x = a.x.x;
-  val.y = a.x.y;
-  b.x = vec_conversion<uint32_t, float2>(val);
-
-  val.x = a.y.x;
-  val.y = a.y.y;
-  b.y = vec_conversion<uint32_t, float2>(val);
-  return b;
-}
-
-// Float4 -> float4
-template <>
-__inline__ __device__ float4 vec_conversion<float4, Float4_>(const Float4_& a) {
-  float4 b;
-  b.x = a.x.x;
-  b.y = a.x.y;
-  b.z = a.y.x;
-  b.w = a.y.y;
-  return b;
-}
-
-// Float8 -> half2x4
-template <>
-__inline__ __device__ uint4 vec_conversion<uint4, Float8_>(const Float8_& a) {
-  uint4 b;
-  b.x = vec_conversion<uint32_t, float2>(a.x);
-  b.y = vec_conversion<uint32_t, float2>(a.y);
-  b.z = vec_conversion<uint32_t, float2>(a.z);
-  b.w = vec_conversion<uint32_t, float2>(a.w);
-  return b;
-}
-
-// float2 -> bfloat162
-template <>
-__inline__ __device__ __nv_bfloat162
-vec_conversion<__nv_bfloat162, float2>(const float2& a) {
-  __nv_bfloat162 b = __float22bfloat162_rn(a);
-  return b;
-}
-
-// Float4 -> bfloat162x2
-template <>
-__inline__ __device__ bf16_4_t
-vec_conversion<bf16_4_t, Float4_>(const Float4_& a) {
-  bf16_4_t b;
-  b.x = __float22bfloat162_rn(a.x);
-  b.y = __float22bfloat162_rn(a.y);
-  return b;
-}
-
-// Float8 -> bfloat162x4
-template <>
-__inline__ __device__ bf16_8_t
-vec_conversion<bf16_8_t, Float8_>(const Float8_& a) {
-  bf16_8_t b;
-  b.x = __float22bfloat162_rn(a.x);
-  b.y = __float22bfloat162_rn(a.y);
-  b.z = __float22bfloat162_rn(a.z);
-  b.w = __float22bfloat162_rn(a.w);
-  return b;
-}
-
 /* Scaled and vectorized conversions, for data exchange between high and low
    precision domains
 
@@ -335,6 +105,7 @@ vec_conversion<bf16_8_t, Float8_>(const Float8_& a) {
  */
 
 using __nv_bfloat16 = __hip_bfloat16;
+using __nv_bfloat162 = __hip_bfloat162;
 
 // fp8 -> __nv_bfloat16
 template <>
@@ -613,17 +384,6 @@ scaled_vec_conversion<uint32_t, float4>(const float4& a, float scale) {
 }
   #endif  // ENABLE_FP8
 
-template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
-__inline__ __device__ Tout convert(const Tin& x) {
-  #ifdef ENABLE_FP8
-  if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) {
-    return vec_conversion<Tout, Tin>(x);
-  }
-  #endif
-  assert(false);
-  return {};  // Squash missing return statement warning
-}
-
 template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
 __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
   #ifdef ENABLE_FP8
diff --git a/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh b/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh
index 11a9b0818d..7fcf541fd6 100644
--- a/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh
+++ b/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh
@@ -282,19 +282,6 @@ __inline__ __device__ float4 scaled_vec_conversion<float4, uint32_t>(
 }
   #endif  // ENABLE_FP8
 
-template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
-__inline__ __device__ Tout convert(const Tin& x) {
-  #if 0  // Disable the following code to reduce the binary size.
-  if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) {
-    return vec_conversion<Tout, Tin>(x, __NV_E4M3);
-  } else if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E5M2) {
-    return vec_conversion<Tout, Tin>(x, __NV_E5M2);
-  }
-  #endif
-  assert(false);
-  __builtin_unreachable();  // Suppress missing return statement warning
-}
-
 template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
 __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
   #ifdef ENABLE_FP8
diff --git a/csrc/rocm/attention.cu b/csrc/rocm/attention.cu
index b747b870e2..2a11653a58 100644
--- a/csrc/rocm/attention.cu
+++ b/csrc/rocm/attention.cu
@@ -342,22 +342,6 @@ typedef union u64_cvt {
   int64_t i64;
 } _T8x8;
 
-__device__ __forceinline__ _B8x8 convert_b16x8(const _B16x8& input,
-                                               _T8x8& Mtemp) {
-  _T8x8 Qtmp8x8;
-
-  for (int i = 0; i < 2; i++) {
-    floatx4 q_out = {0, 0, 0, 0};
-    q_out = gcn_mfma16x16x16_instr<_Float16, 0, 0, 0>(Mtemp.b64, input.xy[i],
-                                                      q_out);
-    Qtmp8x8.b16x4[i * 2] =
-        __builtin_amdgcn_cvt_pk_fp8_f32(q_out[0], q_out[1], 0, false);
-    Qtmp8x8.b16x4[i * 2 + 1] =
-        __builtin_amdgcn_cvt_pk_fp8_f32(q_out[2], q_out[3], 0, false);
-  }
-  return Qtmp8x8.b8x8;
-}
-
 __device__ float warpReduceMax(float val) {
   for (int offset = warpSize / 2; offset > 0; offset /= 2) {
     val = max(
@@ -2583,23 +2567,6 @@ __device__ __forceinline__ float to_float(const T& inp) {
   }
 }
 
-template <typename T>
-__device__ __forceinline__ float to_float_b16(const bit16_t& inp) {
-  union tmpcvt {
-    bit16_t u;
-    _Float16 f;
-    __hip_bfloat16 b;
-  } t16;
-  t16.u = inp;
-  if constexpr (std::is_same<T, _Float16>::value) {
-    return (float)t16.f;
-  } else if constexpr (std::is_same<T, __hip_bfloat16>::value) {
-    return __bfloat162float(t16.b);
-  } else {
-    static_assert(false, "unsupported 16b dtype");
-  }
-}
-
 template <typename T>
 __device__ __forceinline__ T from_float(const float& inp) {
   if constexpr (std::is_same<T, _Float16>::value) {
diff --git a/csrc/rocm/moe_q_gemm_rdna3.cu b/csrc/rocm/moe_q_gemm_rdna3.cu
index 6c25ed7e4b..61ffbb9987 100644
--- a/csrc/rocm/moe_q_gemm_rdna3.cu
+++ b/csrc/rocm/moe_q_gemm_rdna3.cu
@@ -67,21 +67,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
   return result;
 }
 
-__forceinline__ __device__ float dot22_8_f(float (&dq)[8],
-                                           const bf16_t* a_ptr) {
-  float result = 0.0f;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t aw;
-    __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
-    float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
-    float a_y = __uint_as_float(aw & 0xFFFF0000u);
-    result = __fmaf_rn(dq[2 * i + 0], a_x, result);
-    result = __fmaf_rn(dq[2 * i + 1], a_y, result);
-  }
-  return result;
-}
-
 __forceinline__ __device__ void atomic_add_pk4_f16(half* addr, half2 v01,
                                                    half2 v23) {
   unsigned long long* addr_u = reinterpret_cast<unsigned long long*>(addr);
diff --git a/csrc/rocm/q_gemm_rdna3.cu b/csrc/rocm/q_gemm_rdna3.cu
index 104af66d1c..bf23682699 100644
--- a/csrc/rocm/q_gemm_rdna3.cu
+++ b/csrc/rocm/q_gemm_rdna3.cu
@@ -115,61 +115,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
   return result;
 }
 
-__forceinline__ __device__ float dot22_8_f(bf162_t (&dq)[4],
-                                           const bf16_t* a_ptr) {
-  // RDNA3 (gfx1100) lacks a packed bf16 FMA: there is no v_pk_fma_bf16 in
-  // the gfx11 ISA (it only landed on CDNA3+ / gfx94x and later). hipcc
-  // therefore lowers __hfma2(bf162_t, bf162_t, bf162_t) to a serialised
-  // fallback (single-element FMAs or fp32 round-trips), which empirically
-  // runs ~2× the cycle count of v_pk_fma_f16 on the same VALU. The bf16
-  // decode path was paying that tax in full, scaling linearly with M (the
-  // fp16 path scales sub-linearly because its v_pk_fma_f16 is full rate
-  // and the kernel becomes memory-bound).
-  //
-  // Fix: widen bf16 → fp32 explicitly (a left-shift by 16, free in VGPRs)
-  // and accumulate with v_fma_f32, which IS full rate on RDNA3. Same FMA
-  // count, but each FMA is fast. Bonus: the accumulator is now fp32
-  // throughout instead of bf16, which is also numerically more accurate
-  // (no compounding bf16-rounding inside the dot loop).
-  float result = 0.0f;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t aw, dw;
-    __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
-    __builtin_memcpy(&dw, &dq[i], sizeof(uint32_t));
-    // bf16 in low 16 bits  → fp32 by left-shifting into the upper half.
-    // bf16 in high 16 bits → already aligned with fp32's upper half.
-    float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
-    float a_y = __uint_as_float(aw & 0xFFFF0000u);
-    float d_x = __uint_as_float((dw & 0xFFFFu) << 16);
-    float d_y = __uint_as_float(dw & 0xFFFF0000u);
-    result = __fmaf_rn(d_x, a_x, result);
-    result = __fmaf_rn(d_y, a_y, result);
-  }
-  return result;
-}
-
-// fp32-input dot product: paired with dequant_4bit_8_bf16_f32 which already
-// produces fp32 dq[8]. Saves the bf16→fp32 widening that the bf162_t
-// overload above does for dq (still need to widen A from bf16). Wins more
-// at high N: the bf162_t version's per-call widening cost scales with the
-// number of dequants × M_COUNT × 4 dot calls; the fp32 version pays only
-// for A widening (M_COUNT × 4 × 4 widens, half as many).
-__forceinline__ __device__ float dot22_8_f(float (&dq)[8],
-                                           const bf16_t* a_ptr) {
-  float result = 0.0f;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t aw;
-    __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
-    float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
-    float a_y = __uint_as_float(aw & 0xFFFF0000u);
-    result = __fmaf_rn(dq[2 * i + 0], a_x, result);
-    result = __fmaf_rn(dq[2 * i + 1], a_y, result);
-  }
-  return result;
-}
-
 // ---------------------------------------------------------------------------
 // Packed atomic-add via CAS-loop on a 64-bit word (4 fp16/bf16 lanes per CAS).
 // RDNA3 (gfx11) does NOT have native v_global_atomic_pk_add_f16 / _bf16 (those
@@ -327,9 +272,8 @@ __global__ void gemm_q4_kernel_rdna3(const T* __restrict__ a,
 
   // Per-column dequant constants. We hold one set of (z, y) pairs per column.
   // fp16 uses the exllama (z1z16, y1y16) double-pair to enable the upper-
-  // nibble-*16 trick. bf16 uses fp32 scalars (z, y) because the dequant
-  // produces fp32 directly — see prep_zero_scale_bf16_f32 / the FMA
-  // bypass for the missing v_pk_fma_bf16 on gfx11.
+  // nibble-*16 trick. bf16 uses fp32 scalars (z, y) for the inline v_dot2
+  // bias correction.
   half2 z1z16_h[4][2], y1y16_h[4][2];
   float z_b_f[4], y_b_f[4];
 
diff --git a/csrc/rocm/qdq_4_rdna3.cuh b/csrc/rocm/qdq_4_rdna3.cuh
index f4be668c1f..43a7b33ba1 100644
--- a/csrc/rocm/qdq_4_rdna3.cuh
+++ b/csrc/rocm/qdq_4_rdna3.cuh
@@ -113,63 +113,6 @@ __forceinline__ __device__ void dequant_4bit_8_fp16(uint32_t qa, half2 (&dq)[4],
 // bf16 path
 // ---------------------------------------------------------------------------
 
-// Bit-trick magic for bf16:
-//   bf16(128) == 0x4300 (sign 0, exp 134, mantissa 0).
-//   For nibble n in [0..15], bits [3:0] of mantissa hold n exactly because
-//   bf16's ULP at 128 is 1 (mantissa step = 2^(7-7) = 1). So
-//   ((qa & 0x000F000F) | 0x43004300) bitcasts to bfloat162(128+n_lo, 128+n_hi).
-//
-// Because bf16's mantissa is only 7 bits, we cannot use the fp16 "upper nibble
-// * 16" trick. Instead each pair of nibbles is shifted down to [3:0]/[19:16]
-// via a single 4/8/12-bit right-shift before the OR. That costs one extra
-// shift per pair vs fp16, but keeps the FMA structure identical.
-__forceinline__ __device__ void prep_zero_scale_bf16(uint32_t zero,
-                                                     bf16_t scale,
-                                                     bf162_t& z_prep,
-                                                     bf162_t& y_prep) {
-  // z = scale * -(128 + zero); y = scale.
-  float scale_f = __bfloat162float(scale);
-  float zf = -(128.0f + (float)zero) * scale_f;
-  bf16_t zb = __float2bfloat16(zf);
-  z_prep = __bfloat162bfloat162(zb);
-  y_prep = __bfloat162bfloat162(scale);
-}
-
-__forceinline__ __device__ void dequant_4bit_8_bf16(uint32_t qa,
-                                                    bf162_t (&dq)[4],
-                                                    bf162_t z_prep,
-                                                    bf162_t y_prep) {
-  const uint32_t c0 = 0x43004300;
-
-  union {
-    uint32_t u;
-    bf162_t b2;
-  } q0, q1, q2, q3;
-  q0.u = ((qa >> 0) & 0x000F000F) | c0;   // bf162(128+q[0], 128+q[1])
-  q1.u = ((qa >> 4) & 0x000F000F) | c0;   // bf162(128+q[2], 128+q[3])
-  q2.u = ((qa >> 8) & 0x000F000F) | c0;   // bf162(128+q[4], 128+q[5])
-  q3.u = ((qa >> 12) & 0x000F000F) | c0;  // bf162(128+q[6], 128+q[7])
-
-  // dq = q_b * scale + (-(128+zero)*scale) = (q - zero) * scale
-  dq[0] = __hfma2(q0.b2, y_prep, z_prep);
-  dq[1] = __hfma2(q1.b2, y_prep, z_prep);
-  dq[2] = __hfma2(q2.b2, y_prep, z_prep);
-  dq[3] = __hfma2(q3.b2, y_prep, z_prep);
-}
-
-// ---------------------------------------------------------------------------
-// bf16-input → fp32-output dequant (RDNA3 scalar path).
-//
-// RDNA3 (gfx1100) has no v_pk_fma_bf16; packed bf16 FMA lowers to a slow
-// fallback. Rather than computing dq in bf16 and widening at FMA time in
-// the dot product, we widen to fp32 here once (a free left-shift by 16) and
-// emit the (q - zero) * scale FMA directly in fp32. This:
-//   * Replaces 4× slow bf16 packed FMA with 8× fast fp32 FMA per int32.
-//   * Eliminates 4× bf16→fp32 widens that the dot product would do.
-//   * Keeps the dot product accumulator in fp32 without a roundtrip.
-//
-// Output: fp32 dq[8], one element per K position (consumed by the
-// fp32-overload of dot22_8_f in q_gemm_rdna3.cu).
 __forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero,
                                                          bf16_t scale,
                                                          float& z_prep,
@@ -179,60 +122,6 @@ __forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero,
   y_prep = scale_f;
 }
 
-// Pure-q dequant for the M_COUNT=1 factored path: outputs the unscaled fp32
-// values 128+nibble, without folding scale/zero. The caller folds scale/zb
-// into the accumulator outside the inner loop using a precomputed sum_a,
-// which saves ~27% of the FMA count vs the per-col-dequant approach above
-// (only beneficial at M_COUNT=1; break-even at M_COUNT=2).
-//
-// Cost: 0 FMAs (pure bit-trick + as_float reinterprets).
-__forceinline__ __device__ void dequant_4bit_8_bf16_q_only(uint32_t qa,
-                                                           float (&q_f32)[8]) {
-  const uint32_t c0 = 0x43004300;
-  const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
-  const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
-  const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
-  const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
-  q_f32[0] = __uint_as_float((q0 & 0xFFFFu) << 16);
-  q_f32[1] = __uint_as_float(q0 & 0xFFFF0000u);
-  q_f32[2] = __uint_as_float((q1 & 0xFFFFu) << 16);
-  q_f32[3] = __uint_as_float(q1 & 0xFFFF0000u);
-  q_f32[4] = __uint_as_float((q2 & 0xFFFFu) << 16);
-  q_f32[5] = __uint_as_float(q2 & 0xFFFF0000u);
-  q_f32[6] = __uint_as_float((q3 & 0xFFFFu) << 16);
-  q_f32[7] = __uint_as_float(q3 & 0xFFFF0000u);
-}
-
-__forceinline__ __device__ void dequant_4bit_8_bf16_f32(uint32_t qa,
-                                                        float (&dq)[8],
-                                                        float z_prep,
-                                                        float y_prep) {
-  const uint32_t c0 = 0x43004300;
-  const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
-  const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
-  const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
-  const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
-  // bf16(128+nibble) bits → fp32(128+nibble) bits via left-shift by 16
-  // (just zero-extends the mantissa from 7 to 23 bits; exponent preserved).
-  const float q0x = __uint_as_float((q0 & 0xFFFFu) << 16);
-  const float q0y = __uint_as_float(q0 & 0xFFFF0000u);
-  const float q1x = __uint_as_float((q1 & 0xFFFFu) << 16);
-  const float q1y = __uint_as_float(q1 & 0xFFFF0000u);
-  const float q2x = __uint_as_float((q2 & 0xFFFFu) << 16);
-  const float q2y = __uint_as_float(q2 & 0xFFFF0000u);
-  const float q3x = __uint_as_float((q3 & 0xFFFFu) << 16);
-  const float q3y = __uint_as_float(q3 & 0xFFFF0000u);
-  // dq[i] = q_f32 * scale + (-(128+zero)*scale) = (nibble - zero) * scale
-  dq[0] = __fmaf_rn(q0x, y_prep, z_prep);
-  dq[1] = __fmaf_rn(q0y, y_prep, z_prep);
-  dq[2] = __fmaf_rn(q1x, y_prep, z_prep);
-  dq[3] = __fmaf_rn(q1y, y_prep, z_prep);
-  dq[4] = __fmaf_rn(q2x, y_prep, z_prep);
-  dq[5] = __fmaf_rn(q2y, y_prep, z_prep);
-  dq[6] = __fmaf_rn(q3x, y_prep, z_prep);
-  dq[7] = __fmaf_rn(q3y, y_prep, z_prep);
-}
-
 }  // namespace gptq_rdna3
 }  // namespace vllm