diff --git a/aphrodite/common/envs.py b/aphrodite/common/envs.py
index a67c62c9ba..2e68c6c9a3 100755
--- a/aphrodite/common/envs.py
+++ b/aphrodite/common/envs.py
@@ -86,6 +86,7 @@
     APHRODITE_USE_TRITON_AWQ: bool = False
     APHRODITE_ALLOW_RUNTIME_LORA_UPDATING: bool = False
     APHRODITE_SKIP_P2P_CHECK: bool = False
+    APHRODITE_FORCE_P2P: bool = False
     APHRODITE_DISABLED_KERNELS: list[str] = []
     APHRODITE_USE_V1: bool = True
     APHRODITE_ROCM_USE_AITER: bool = False
@@ -717,6 +718,11 @@ def get_aphrodite_port() -> Optional[int]:
     "APHRODITE_SKIP_P2P_CHECK":
     lambda: os.getenv("APHRODITE_SKIP_P2P_CHECK", "1") == "1",
 
+    # If set, Aphrodite will skip the P2P check and assume that P2P is
+    # available. Used for custom all-reduce kernels.
+    "APHRODITE_FORCE_P2P":
+    lambda: bool(int(os.getenv("APHRODITE_FORCE_P2P", "0"))),
+
     # List of quantization kernels that should be disabled, used for testing
     # and performance comparisons. Currently only affects MPLinearKernel
     # selection
diff --git a/aphrodite/distributed/device_communicators/custom_all_reduce.py b/aphrodite/distributed/device_communicators/custom_all_reduce.py
index 9572020840..f260548d49 100644
--- a/aphrodite/distributed/device_communicators/custom_all_reduce.py
+++ b/aphrodite/distributed/device_communicators/custom_all_reduce.py
@@ -141,13 +141,14 @@ def __init__(self,
         assert current_platform.is_cuda_alike()
         fully_connected = current_platform.is_fully_connected(
             physical_device_ids)
-        if world_size > 2 and not fully_connected:
-            if self.rank == 0:
-                logger.warning(
-                    "Custom allreduce is disabled because it's not supported on"
-                    " more than two PCIe-only GPUs. To silence this warning, "
-                    "specify disable_custom_all_reduce=True explicitly.")
-            return
+        if not envs.APHRODITE_FORCE_P2P:
+            if world_size > 2 and not fully_connected:
+                if self.rank == 0:
+                    logger.warning(
+                        "Custom allreduce is disabled because it's not supported on"
+                        " more than two PCIe-only GPUs. To silence this warning, "
+                        "specify disable_custom_all_reduce=True explicitly.")
+                return
         # test P2P capability, this checks software/cudaruntime support
         # this is expensive to compute at the first time
         # then we cache the result
diff --git a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4.cu b/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4.cu
deleted file mode 100644
index b23cf1d921..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4.cu
+++ /dev/null
@@ -1,89 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}
diff --git a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4b8.cu b/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4b8.cu
deleted file mode 100644
index d6602efbff..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku4b8.cu
+++ /dev/null
@@ -1,109 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU4B8.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}
diff --git a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku8b128.cu b/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku8b128.cu
deleted file mode 100644
index 121e784b01..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_bf16_ite::ku8b128.cu
+++ /dev/null
@@ -1,109 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<nv_bfloat16, aphrodite::kU8B128.id(), aphrodite::kBFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}
diff --git a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4.cu b/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4.cu
deleted file mode 100644
index 3e12c2016a..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4.cu
+++ /dev/null
@@ -1,89 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}
diff --git a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4b8.cu b/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4b8.cu
deleted file mode 100644
index 46e62b623b..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku4b8.cu
+++ /dev/null
@@ -1,109 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU4B8.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}
diff --git a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku8b128.cu b/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku8b128.cu
deleted file mode 100644
index e9b524eebe..0000000000
--- a/kernels/moe/marlin_moe_wna16/kernel_fp16_ite::ku8b128.cu
+++ /dev/null
@@ -1,109 +0,0 @@
-// auto generated by generate.py
-// clang-format off
-
-#include "kernel.h"
-#include "marlin_template.h"
-
-namespace MARLIN_NAMESPACE_NAME {
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 0, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, -1, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 2, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 4, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, true, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 1, 8, 8, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 1, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 2, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 2, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 3, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 3, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 256, 4, 16, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-template __global__ void Marlin<half, aphrodite::kU8B128.id(), aphrodite::kFloat16.id(), 128, 4, 8, 4, false, pipe_stages, 8, false>( MARLIN_KERNEL_PARAMS );
-
-}