diff --git a/ex_engine/xllm_kernels/cuda/block_copy.cu b/ex_engine/xllm_kernels/cuda/block_copy.cu
index 23ef742b..d92e7b3e 100644
--- a/ex_engine/xllm_kernels/cuda/block_copy.cu
+++ b/ex_engine/xllm_kernels/cuda/block_copy.cu
@@ -25,10 +25,6 @@ limitations under the License.
 
 #include "device_utils.cuh"
 
-#ifndef DEVICE_INLINE
-#define DEVICE_INLINE __device__ __forceinline__
-#define HOST_DEVICE_INLINE __host__ __device__ __forceinline__
-#endif
 
 
 namespace xllm::kernel::cuda {
diff --git a/ex_engine/xllm_kernels/cuda/headers/device_utils.cuh b/ex_engine/xllm_kernels/cuda/headers/device_utils.cuh
index d14a6cb5..7115fcfc 100644
--- a/ex_engine/xllm_kernels/cuda/headers/device_utils.cuh
+++ b/ex_engine/xllm_kernels/cuda/headers/device_utils.cuh
@@ -115,6 +115,20 @@ struct TopkConstants {
 
 }  // namespace xllm::kernel::cuda
 
+// ============================================================================
+// Portable macros and utilities (from xllm/core/kernels/cuda/utils.h)
+// ============================================================================
+#ifndef DEVICE_INLINE
+#define DEVICE_INLINE __device__ __forceinline__
+#define HOST_DEVICE_INLINE __host__ __device__ __forceinline__
+#endif
+
+template <typename T>
+HOST_DEVICE_INLINE constexpr std::enable_if_t<std::is_integral_v<T>, T>
+ceil_div(T a, T b) {
+    return (a + b - 1) / b;
+}
+
 // ============================================================================
 // Dispatch macros (from xllm/core/kernels/cuda/utils.h)
 // These wrap AT_DISPATCH_SWITCH for float16/bfloat16/float32 dispatch.