diff --git a/ex_engine/xllm_kernels/cuda/activation.cu b/ex_engine/xllm_kernels/cuda/activation.cu
index 1f4a71e5..a7cafc63 100644
--- a/ex_engine/xllm_kernels/cuda/activation.cu
+++ b/ex_engine/xllm_kernels/cuda/activation.cu
@@ -17,7 +17,7 @@ limitations under the License.
 
 #include <cstdint>
 
-#include "cuda_ops_api.h"
+
 #include "device_utils.cuh"
 
 // ref to:
diff --git a/ex_engine/xllm_kernels/cuda/block_copy.cu b/ex_engine/xllm_kernels/cuda/block_copy.cu
index 05651058..0b46c325 100644
--- a/ex_engine/xllm_kernels/cuda/block_copy.cu
+++ b/ex_engine/xllm_kernels/cuda/block_copy.cu
@@ -23,7 +23,7 @@ limitations under the License.
 #include <cstdint>
 #include <type_traits>
 
-#include "cuda_ops_api.h"
+#include "device_utils.cuh"
 #include "utils.h"
 
 namespace xllm::kernel::cuda {
diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_combine.cu b/ex_engine/xllm_kernels/cuda/moe/moe_combine.cu
index f4f21c69..4be4078c 100755
--- a/ex_engine/xllm_kernels/cuda/moe/moe_combine.cu
+++ b/ex_engine/xllm_kernels/cuda/moe/moe_combine.cu
@@ -28,7 +28,7 @@ limitations under the License.
 #include <c10/cuda/CUDAGuard.h>
 
 #include "device_utils.cuh"
-#include "kernels/cuda/cuda_ops_api.h"
+#include "device_utils.cuh"
 
 namespace xllm::kernel::cuda {
 
diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_compute_index.cu b/ex_engine/xllm_kernels/cuda/moe/moe_compute_index.cu
index 5e15a442..b7eae3bf 100644
--- a/ex_engine/xllm_kernels/cuda/moe/moe_compute_index.cu
+++ b/ex_engine/xllm_kernels/cuda/moe/moe_compute_index.cu
@@ -27,7 +27,7 @@ limitations under the License.
 
 #include <cub/block/block_scan.cuh>
 
-#include "kernels/cuda/cuda_ops_api.h"
+#include "device_utils.cuh"
 
 namespace xllm::kernel::cuda {
 
diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_fused_topk.cu b/ex_engine/xllm_kernels/cuda/moe/moe_fused_topk.cu
index 9aaf0c12..186be86e 100644
--- a/ex_engine/xllm_kernels/cuda/moe/moe_fused_topk.cu
+++ b/ex_engine/xllm_kernels/cuda/moe/moe_fused_topk.cu
@@ -15,7 +15,7 @@ limitations under the License.
 #if defined(USE_DCU)
 #include "kernels/dcu/dcu_ops_api.h"
 #else
-#include "kernels/cuda/cuda_ops_api.h"
+#include "device_utils.cuh"
 #endif
 #include "moe_topk_sigmoid_kernels.cuh"
 #include "moe_topk_softmax_kernels.cuh"
diff --git a/ex_engine/xllm_kernels/cuda/norm.cu b/ex_engine/xllm_kernels/cuda/norm.cu
index 06c4a8c7..511ca66b 100644
--- a/ex_engine/xllm_kernels/cuda/norm.cu
+++ b/ex_engine/xllm_kernels/cuda/norm.cu
@@ -19,7 +19,6 @@ limitations under the License.
 #include <cstdint>
 #include <cub/cub.cuh>
 
-#include "cuda_ops_api.h"
 #include "device_utils.cuh"
 #include "fp8_quant_utils.cuh"
 #include "type_convert.cuh"
@@ -27,14 +26,10 @@ limitations under the License.
 // ref to:
 // https://github.com/vllm-project/vllm/blob/main/csrc/layernorm_kernels.cu
 
-#if CUB_VERSION >= 200800
-// corex CUB (CUDA 10.2) — no cuda::std::functional
+// corex CUB (CUDA 10.2) — use old-style CUB operators
 using CubAddOp = cub::Sum;
 using CubMaxOp = cub::Max;
-#else   // if CUB_VERSION < 200800
-using CubAddOp = cub::Sum;
-using CubMaxOp = cub::Max;
-#endif  // CUB_VERSION
+
 
 namespace {
 
diff --git a/ex_engine/xllm_kernels/cuda/reshape_paged_cache.cu b/ex_engine/xllm_kernels/cuda/reshape_paged_cache.cu
index 3fe3db91..c835f95c 100644
--- a/ex_engine/xllm_kernels/cuda/reshape_paged_cache.cu
+++ b/ex_engine/xllm_kernels/cuda/reshape_paged_cache.cu
@@ -15,7 +15,7 @@ limitations under the License.
 
 #include <c10/cuda/CUDAStream.h>
 
-#include "cuda_ops_api.h"
+
 #include "device_utils.cuh"
 
 namespace xllm::kernel::cuda {
diff --git a/ex_engine/xllm_kernels/cuda/rope.cu b/ex_engine/xllm_kernels/cuda/rope.cu
index 585cdf99..856207f2 100644
--- a/ex_engine/xllm_kernels/cuda/rope.cu
+++ b/ex_engine/xllm_kernels/cuda/rope.cu
@@ -16,7 +16,7 @@ limitations under the License.
 #include <c10/cuda/CUDAGuard.h>
 #include <torch/all.h>
 
-#include "cuda_ops_api.h"
+
 #include "device_utils.cuh"
 
 // ref to: