diff --git a/src/native/cuda/mars/blas.h b/src/native/cuda/mars/blas.h
new file mode 100644
index 000000000..25ced1dd0
--- /dev/null
+++ b/src/native/cuda/mars/blas.h
@@ -0,0 +1,51 @@
+#ifndef INFINI_OPS_MARS_BLAS_H_
+#define INFINI_OPS_MARS_BLAS_H_
+
+#include <utility>
+
+// clang-format off
+#include <hcblas/hcblas.h>
+// clang-format on
+
+#include "data_type.h"
+#include "native/cuda/blas.h"
+#include "native/cuda/mars/blas_utils.h"
+#include "native/cuda/mars/runtime_.h"
+
+namespace infini::ops {
+
+template <>
+struct Blas<Device::Type::kMars> : public Runtime<Device::Type::kMars> {
+  using BlasHandle = hcblasHandle_t;
+
+  static constexpr auto BLAS_OP_N = HCBLAS_OP_N;
+
+  static constexpr auto BLAS_OP_T = HCBLAS_OP_T;
+
+  static constexpr auto R_16F = HPCC_R_16F;
+
+  static constexpr auto R_16BF = HPCC_R_16BF;
+
+  static constexpr auto R_32F = HPCC_R_32F;
+
+  static constexpr auto BLAS_COMPUTE_32F = HCBLAS_COMPUTE_32F;
+
+  static constexpr auto BLAS_COMPUTE_32F_FAST_TF32 =
+      HCBLAS_COMPUTE_32F_FAST_TF32;
+
+  static constexpr auto BLAS_GEMM_DEFAULT = HCBLAS_GEMM_DEFAULT;
+
+  static constexpr auto BlasCreate = hcblasCreate;
+
+  static constexpr auto BlasSetStream = hcblasSetStream;
+
+  static constexpr auto BlasDestroy = hcblasDestroy;
+
+  static constexpr auto BlasGemmStridedBatchedEx = [](auto&&... args) {
+    return hcblasGemmStridedBatchedEx(std::forward<decltype(args)>(args)...);
+  };
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/blas_utils.h b/src/native/cuda/mars/blas_utils.h
new file mode 100644
index 000000000..b6e4bc38c
--- /dev/null
+++ b/src/native/cuda/mars/blas_utils.h
@@ -0,0 +1,30 @@
+#ifndef INFINI_OPS_MARS_BLAS_UTILS_H_
+#define INFINI_OPS_MARS_BLAS_UTILS_H_
+
+// clang-format off
+#include <hcblas/hcblas.h>
+// clang-format on
+
+#include "data_type.h"
+#include "native/cuda/blas_utils.h"
+
+namespace infini::ops {
+
+template <>
+struct BlasUtils<Device::Type::kMars> {
+  static auto GetDataType(DataType dtype) {
+    if (dtype == DataType::kFloat16) return HPCC_R_16F;
+    if (dtype == DataType::kBFloat16) return HPCC_R_16BF;
+    return HPCC_R_32F;
+  }
+
+  static auto GetComputeType(DataType dtype) {
+    if (dtype == DataType::kFloat16 || dtype == DataType::kBFloat16)
+      return HCBLAS_COMPUTE_32F;
+    return HCBLAS_COMPUTE_32F_FAST_TF32;
+  }
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/caster.cuh b/src/native/cuda/mars/caster.cuh
new file mode 100644
index 000000000..6d51a151a
--- /dev/null
+++ b/src/native/cuda/mars/caster.cuh
@@ -0,0 +1,78 @@
+#ifndef INFINI_OPS_MARS_CASTER__H_
+#define INFINI_OPS_MARS_CASTER__H_
+
+#include "native/cuda/caster.cuh"
+#include "native/cuda/mars/data_type_.h"
+
+namespace infini::ops {
+
+namespace detail {
+
+template <>
+struct ToFloat<Device::Type::kMars, __half> {
+  __host__ __device__ float operator()(__half x) { return __half2float(x); }
+};
+
+template <>
+struct ToFloat<Device::Type::kMars, __hpcc_bfloat16> {
+  __host__ __device__ float operator()(__hpcc_bfloat16 x) {
+    return __bfloat162float(x);
+  }
+};
+
+template <>
+struct FromFloat<Device::Type::kMars, __half> {
+  __host__ __device__ __half operator()(float f) { return __float2half(f); }
+};
+
+template <>
+struct FromFloat<Device::Type::kMars, __hpcc_bfloat16> {
+  __host__ __device__ __hpcc_bfloat16 operator()(float f) {
+    return __float2bfloat16(f);
+  }
+};
+
+template <>
+struct HardwareCast<Device::Type::kMars, __hpcc_bfloat16, int> {
+  inline static constexpr bool kSupported = true;
+  __host__ __device__ __hpcc_bfloat16 operator()(int x) {
+    return __int2bfloat16_rn(x);
+  }
+};
+
+template <>
+struct HardwareCast<Device::Type::kMars, __half, int> {
+  inline static constexpr bool kSupported = true;
+  __host__ __device__ __half operator()(int x) { return __int2half_rn(x); }
+};
+
+template <>
+struct HardwareCast<Device::Type::kMars, __hpcc_bfloat16, double> {
+  inline static constexpr bool kSupported = true;
+  __host__ __device__ __hpcc_bfloat16 operator()(double x) {
+    return __double2bfloat16(x);
+  }
+};
+
+template <>
+struct HardwareCast<Device::Type::kMars, __half, double> {
+  inline static constexpr bool kSupported = true;
+  __host__ __device__ __half operator()(double x) { return __double2half(x); }
+};
+
+template <>
+struct HardwareCast<Device::Type::kMars, __half, __hpcc_bfloat16> {
+  inline static constexpr bool kSupported = true;
+  __host__ __device__ __half operator()(__hpcc_bfloat16 x) {
+    return __float2half_rn(__bfloat162float(x));
+  }
+};
+
+}  // namespace detail
+
+template <>
+struct Caster<Device::Type::kMars> : CudaCasterImpl<Device::Type::kMars> {};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/data_type_.h b/src/native/cuda/mars/data_type_.h
new file mode 100644
index 000000000..140e300c8
--- /dev/null
+++ b/src/native/cuda/mars/data_type_.h
@@ -0,0 +1,13 @@
+#ifndef INFINI_OPS_MARS_DATA_TYPE__H_
+#define INFINI_OPS_MARS_DATA_TYPE__H_
+
+#include <infini/rt/mars/data_type_.h>
+
+namespace infini::ops {
+
+using infini::rt::cuda_bfloat16;
+using infini::rt::cuda_bfloat162;
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/device_.h b/src/native/cuda/mars/device_.h
new file mode 100644
index 000000000..eb7158d1a
--- /dev/null
+++ b/src/native/cuda/mars/device_.h
@@ -0,0 +1,6 @@
+#ifndef INFINI_OPS_MARS_DEVICE__H_
+#define INFINI_OPS_MARS_DEVICE__H_
+
+#include <infini/rt/mars/device_.h>
+
+#endif
diff --git a/src/native/cuda/mars/device_property.h b/src/native/cuda/mars/device_property.h
new file mode 100644
index 000000000..0700a39a6
--- /dev/null
+++ b/src/native/cuda/mars/device_property.h
@@ -0,0 +1,11 @@
+#ifndef INFINI_OPS_MARS_DEVICE_PROPERTY_H_
+#define INFINI_OPS_MARS_DEVICE_PROPERTY_H_
+
+namespace infini::ops {
+
+// TODO: Add HCR device properties query for Mars.
+inline int QueryMaxThreadsPerBlock() { return 256; }
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/add/kernel.h b/src/native/cuda/mars/ops/add/kernel.h
new file mode 100644
index 000000000..f94613cda
--- /dev/null
+++ b/src/native/cuda/mars/ops/add/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_ADD_KERNEL_H_
+#define INFINI_OPS_MARS_ADD_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/add/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Add, Device::Type::kMars>
+    : public CudaAdd<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaAdd<Runtime<Device::Type::kMars>>::CudaAdd;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/add_rms_norm/kernel.h b/src/native/cuda/mars/ops/add_rms_norm/kernel.h
new file mode 100644
index 000000000..5d7bf0b6a
--- /dev/null
+++ b/src/native/cuda/mars/ops/add_rms_norm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_ADD_RMS_NORM_KERNEL_H_
+#define INFINI_OPS_MARS_ADD_RMS_NORM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/add_rms_norm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<AddRmsNorm, Device::Type::kMars>
+    : public CudaAddRmsNorm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaAddRmsNorm<Runtime<Device::Type::kMars>>::CudaAddRmsNorm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/causal_softmax/kernel.h b/src/native/cuda/mars/ops/causal_softmax/kernel.h
new file mode 100644
index 000000000..df3b5579e
--- /dev/null
+++ b/src/native/cuda/mars/ops/causal_softmax/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CAUSAL_SOFTMAX_KERNEL_H_
+#define INFINI_OPS_MARS_CAUSAL_SOFTMAX_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/causal_softmax/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<CausalSoftmax, Device::Type::kMars>
+    : public CudaCausalSoftmax<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaCausalSoftmax<Runtime<Device::Type::kMars>>::CudaCausalSoftmax;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h b/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h
new file mode 100644
index 000000000..f97239ffb
--- /dev/null
+++ b/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_CAUSAL_SOFTMAX_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_CAUSAL_SOFTMAX_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/causal_softmax_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<CausalSoftmaxInfinilm, Device::Type::kMars>
+    : public CudaCausalSoftmaxInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaCausalSoftmaxInfinilm<
+      Runtime<Device::Type::kMars>>::CudaCausalSoftmaxInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/conv1d/kernel.h b/src/native/cuda/mars/ops/conv1d/kernel.h
new file mode 100644
index 000000000..f855b7fea
--- /dev/null
+++ b/src/native/cuda/mars/ops/conv1d/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CONV1D_KERNEL_H_
+#define INFINI_OPS_MARS_CONV1D_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/convolution/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Conv1d, Device::Type::kMars>
+    : public CudaConv<Runtime<Device::Type::kMars>, Conv1d> {
+ public:
+  using CudaConv<Runtime<Device::Type::kMars>, Conv1d>::CudaConv;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/conv2d/kernel.h b/src/native/cuda/mars/ops/conv2d/kernel.h
new file mode 100644
index 000000000..e3b9ef5c0
--- /dev/null
+++ b/src/native/cuda/mars/ops/conv2d/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CONV2D_KERNEL_H_
+#define INFINI_OPS_MARS_CONV2D_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/convolution/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Conv2d, Device::Type::kMars>
+    : public CudaConv<Runtime<Device::Type::kMars>, Conv2d> {
+ public:
+  using CudaConv<Runtime<Device::Type::kMars>, Conv2d>::CudaConv;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/conv3d/kernel.h b/src/native/cuda/mars/ops/conv3d/kernel.h
new file mode 100644
index 000000000..02ff7b5cf
--- /dev/null
+++ b/src/native/cuda/mars/ops/conv3d/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CONV3D_KERNEL_H_
+#define INFINI_OPS_MARS_CONV3D_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/convolution/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Conv3d, Device::Type::kMars>
+    : public CudaConv<Runtime<Device::Type::kMars>, Conv3d> {
+ public:
+  using CudaConv<Runtime<Device::Type::kMars>, Conv3d>::CudaConv;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/conv_infinilm/kernel.h b/src/native/cuda/mars/ops/conv_infinilm/kernel.h
new file mode 100644
index 000000000..f5010866f
--- /dev/null
+++ b/src/native/cuda/mars/ops/conv_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CONV_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_CONV_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/conv_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<ConvInfinilm, Device::Type::kMars>
+    : public CudaConvInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaConvInfinilm<Runtime<Device::Type::kMars>>::CudaConvInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/convolution/kernel.h b/src/native/cuda/mars/ops/convolution/kernel.h
new file mode 100644
index 000000000..a69289ffc
--- /dev/null
+++ b/src/native/cuda/mars/ops/convolution/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_CONVOLUTION_KERNEL_H_
+#define INFINI_OPS_MARS_CONVOLUTION_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/convolution/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Convolution, Device::Type::kMars>
+    : public CudaConv<Runtime<Device::Type::kMars>, Convolution> {
+ public:
+  using CudaConv<Runtime<Device::Type::kMars>, Convolution>::CudaConv;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/copy/kernel.h b/src/native/cuda/mars/ops/copy/kernel.h
new file mode 100644
index 000000000..36653d910
--- /dev/null
+++ b/src/native/cuda/mars/ops/copy/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_COPY_KERNEL_H_
+#define INFINI_OPS_MARS_COPY_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/copy/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Copy, Device::Type::kMars>
+    : public CudaCopy<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaCopy<Runtime<Device::Type::kMars>>::CudaCopy;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/embedding/kernel.h b/src/native/cuda/mars/ops/embedding/kernel.h
new file mode 100644
index 000000000..f98ac3da7
--- /dev/null
+++ b/src/native/cuda/mars/ops/embedding/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_EMBEDDING_KERNEL_H_
+#define INFINI_OPS_MARS_EMBEDDING_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/embedding/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Embedding, Device::Type::kMars>
+    : public CudaEmbedding<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaEmbedding<Runtime<Device::Type::kMars>>::CudaEmbedding;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/fill/kernel.h b/src/native/cuda/mars/ops/fill/kernel.h
new file mode 100644
index 000000000..90d40c485
--- /dev/null
+++ b/src/native/cuda/mars/ops/fill/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_FILL_KERNEL_H_
+#define INFINI_OPS_MARS_FILL_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/fill/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Fill, Device::Type::kMars>
+    : public CudaFill<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaFill<Runtime<Device::Type::kMars>>::CudaFill;
+};
+
+}  // namespace infini::ops
+
+#endif  // INFINI_OPS_MARS_FILL_KERNEL_H_
diff --git a/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h b/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h
new file mode 100644
index 000000000..666d94474
--- /dev/null
+++ b/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_FUSED_ADD_RMS_NORM_KERNEL_H_
+#define INFINI_OPS_MARS_FUSED_ADD_RMS_NORM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/fused_add_rms_norm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<FusedAddRmsNorm, Device::Type::kMars>
+    : public CudaFusedAddRmsNorm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaFusedAddRmsNorm<Runtime<Device::Type::kMars>>::CudaFusedAddRmsNorm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/gelu/kernel.h b/src/native/cuda/mars/ops/gelu/kernel.h
new file mode 100644
index 000000000..67b88702d
--- /dev/null
+++ b/src/native/cuda/mars/ops/gelu/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_GELU_KERNEL_H_
+#define INFINI_OPS_MARS_GELU_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/gelu/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Gelu, Device::Type::kMars>
+    : public CudaGelu<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaGelu<Runtime<Device::Type::kMars>>::CudaGelu;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/gelu_infinilm/kernel.h b/src/native/cuda/mars/ops/gelu_infinilm/kernel.h
new file mode 100644
index 000000000..6272fb8b8
--- /dev/null
+++ b/src/native/cuda/mars/ops/gelu_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_GELU_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_GELU_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/gelu_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<GeluInfinilm, Device::Type::kMars>
+    : public CudaGeluInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaGeluInfinilm<Runtime<Device::Type::kMars>>::CudaGeluInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h b/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h
new file mode 100644
index 000000000..9ae8bce99
--- /dev/null
+++ b/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_GELUTANH_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_GELUTANH_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/gelutanh_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<GelutanhInfinilm, Device::Type::kMars>
+    : public CudaGelutanhInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaGelutanhInfinilm<
+      Runtime<Device::Type::kMars>>::CudaGelutanhInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/gemm/hcblas.h b/src/native/cuda/mars/ops/gemm/hcblas.h
new file mode 100644
index 000000000..be00e99de
--- /dev/null
+++ b/src/native/cuda/mars/ops/gemm/hcblas.h
@@ -0,0 +1,18 @@
+#ifndef INFINI_OPS_MARS_GEMM_HCBLAS_H_
+#define INFINI_OPS_MARS_GEMM_HCBLAS_H_
+
+#include "native/cuda/mars/blas.h"
+#include "native/cuda/ops/gemm/blas.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Gemm, Device::Type::kMars>
+    : public BlasGemm<Blas<Device::Type::kMars>> {
+ public:
+  using BlasGemm<Blas<Device::Type::kMars>>::BlasGemm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h b/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h
new file mode 100644
index 000000000..a46addbd9
--- /dev/null
+++ b/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_KV_CACHING_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_KV_CACHING_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/kv_caching_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<KvCachingInfinilm, Device::Type::kMars>
+    : public CudaKvCachingInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaKvCachingInfinilm<
+      Runtime<Device::Type::kMars>>::CudaKvCachingInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/mul/kernel.h b/src/native/cuda/mars/ops/mul/kernel.h
new file mode 100644
index 000000000..d4d2520c3
--- /dev/null
+++ b/src/native/cuda/mars/ops/mul/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_MUL_KERNEL_H_
+#define INFINI_OPS_MARS_MUL_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/mul/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Mul, Device::Type::kMars>
+    : public CudaMul<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaMul<Runtime<Device::Type::kMars>>::CudaMul;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h
new file mode 100644
index 000000000..4ac03add9
--- /dev/null
+++ b/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_PAGED_ATTENTION_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_PAGED_ATTENTION_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/paged_attention_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<PagedAttentionInfinilm, Device::Type::kMars>
+    : public CudaPagedAttentionInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaPagedAttentionInfinilm<
+      Runtime<Device::Type::kMars>>::CudaPagedAttentionInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h
new file mode 100644
index 000000000..3825d1002
--- /dev/null
+++ b/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_PAGED_ATTENTION_PREFILL_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_PAGED_ATTENTION_PREFILL_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/paged_attention_prefill_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<PagedAttentionPrefillInfinilm, Device::Type::kMars>
+    : public CudaPagedAttentionPrefillInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaPagedAttentionPrefillInfinilm<
+      Runtime<Device::Type::kMars>>::CudaPagedAttentionPrefillInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h
new file mode 100644
index 000000000..b95b3e252
--- /dev/null
+++ b/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_PAGED_CACHING_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_PAGED_CACHING_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/paged_caching_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<PagedCachingInfinilm, Device::Type::kMars>
+    : public CudaPagedCachingInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaPagedCachingInfinilm<
+      Runtime<Device::Type::kMars>>::CudaPagedCachingInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h b/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h
new file mode 100644
index 000000000..2861368c1
--- /dev/null
+++ b/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_RANDOM_SAMPLE_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_RANDOM_SAMPLE_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/random_sample_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<RandomSampleInfinilm, Device::Type::kMars>
+    : public CudaRandomSampleInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRandomSampleInfinilm<
+      Runtime<Device::Type::kMars>>::CudaRandomSampleInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h b/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h
new file mode 100644
index 000000000..85b9e235d
--- /dev/null
+++ b/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_REARRANGE_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_REARRANGE_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/rearrange_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<RearrangeInfinilm, Device::Type::kMars>
+    : public CudaRearrangeInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRearrangeInfinilm<
+      Runtime<Device::Type::kMars>>::CudaRearrangeInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/relu/kernel.h b/src/native/cuda/mars/ops/relu/kernel.h
new file mode 100644
index 000000000..8fa6b8c76
--- /dev/null
+++ b/src/native/cuda/mars/ops/relu/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_RELU_KERNEL_H_
+#define INFINI_OPS_MARS_RELU_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/relu/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Relu, Device::Type::kMars>
+    : public CudaRelu<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRelu<Runtime<Device::Type::kMars>>::CudaRelu;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/relu_infinilm/kernel.h b/src/native/cuda/mars/ops/relu_infinilm/kernel.h
new file mode 100644
index 000000000..e7810c07b
--- /dev/null
+++ b/src/native/cuda/mars/ops/relu_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_RELU_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_RELU_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/relu_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<ReluInfinilm, Device::Type::kMars>
+    : public CudaReluInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaReluInfinilm<Runtime<Device::Type::kMars>>::CudaReluInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h b/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h
new file mode 100644
index 000000000..c8a022a77
--- /dev/null
+++ b/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_RESHAPE_AND_CACHE_FLASH_KERNEL_H_
+#define INFINI_OPS_MARS_RESHAPE_AND_CACHE_FLASH_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/reshape_and_cache_flash/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<ReshapeAndCacheFlash, Device::Type::kMars>
+    : public CudaReshapeAndCacheFlash<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaReshapeAndCacheFlash<
+      Runtime<Device::Type::kMars>>::CudaReshapeAndCacheFlash;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/rms_norm/kernel.h b/src/native/cuda/mars/ops/rms_norm/kernel.h
new file mode 100644
index 000000000..3bc9f79e5
--- /dev/null
+++ b/src/native/cuda/mars/ops/rms_norm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_RMS_NORM_KERNEL_H_
+#define INFINI_OPS_MARS_RMS_NORM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/rms_norm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<RmsNorm, Device::Type::kMars>
+    : public CudaRmsNorm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRmsNorm<Runtime<Device::Type::kMars>>::CudaRmsNorm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/rotary_embedding/kernel.h b/src/native/cuda/mars/ops/rotary_embedding/kernel.h
new file mode 100644
index 000000000..fe5d73659
--- /dev/null
+++ b/src/native/cuda/mars/ops/rotary_embedding/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_ROTARY_EMBEDDING_KERNEL_H_
+#define INFINI_OPS_MARS_ROTARY_EMBEDDING_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/rotary_embedding/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<RotaryEmbedding, Device::Type::kMars>
+    : public CudaRotaryEmbedding<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRotaryEmbedding<Runtime<Device::Type::kMars>>::CudaRotaryEmbedding;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h b/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h
new file mode 100644
index 000000000..38629b6ae
--- /dev/null
+++ b/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_ROTARY_EMBEDDING_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_ROTARY_EMBEDDING_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/rotary_embedding_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<RotaryEmbeddingInfinilm, Device::Type::kMars>
+    : public CudaRotaryEmbeddingInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaRotaryEmbeddingInfinilm<
+      Runtime<Device::Type::kMars>>::CudaRotaryEmbeddingInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/sigmoid/kernel.h b/src/native/cuda/mars/ops/sigmoid/kernel.h
new file mode 100644
index 000000000..fbb6367eb
--- /dev/null
+++ b/src/native/cuda/mars/ops/sigmoid/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SIGMOID_KERNEL_H_
+#define INFINI_OPS_MARS_SIGMOID_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/sigmoid/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Sigmoid, Device::Type::kMars>
+    : public CudaSigmoid<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSigmoid<Runtime<Device::Type::kMars>>::CudaSigmoid;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h b/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h
new file mode 100644
index 000000000..cde257122
--- /dev/null
+++ b/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SIGMOID_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_SIGMOID_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/sigmoid_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<SigmoidInfinilm, Device::Type::kMars>
+    : public CudaSigmoidInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSigmoidInfinilm<Runtime<Device::Type::kMars>>::CudaSigmoidInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/silu/kernel.h b/src/native/cuda/mars/ops/silu/kernel.h
new file mode 100644
index 000000000..2336ac20d
--- /dev/null
+++ b/src/native/cuda/mars/ops/silu/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SILU_KERNEL_H_
+#define INFINI_OPS_MARS_SILU_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/silu/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Silu, Device::Type::kMars>
+    : public CudaSilu<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSilu<Runtime<Device::Type::kMars>>::CudaSilu;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/silu_and_mul/kernel.h b/src/native/cuda/mars/ops/silu_and_mul/kernel.h
new file mode 100644
index 000000000..bed125e31
--- /dev/null
+++ b/src/native/cuda/mars/ops/silu_and_mul/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SILU_AND_MUL_KERNEL_H_
+#define INFINI_OPS_MARS_SILU_AND_MUL_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/silu_and_mul/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<SiluAndMul, Device::Type::kMars>
+    : public CudaSiluAndMul<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSiluAndMul<Runtime<Device::Type::kMars>>::CudaSiluAndMul;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h b/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h
new file mode 100644
index 000000000..45e136208
--- /dev/null
+++ b/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_SILU_AND_MUL_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_SILU_AND_MUL_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/silu_and_mul_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<SiluAndMulInfinilm, Device::Type::kMars>
+    : public CudaSiluAndMulInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSiluAndMulInfinilm<
+      Runtime<Device::Type::kMars>>::CudaSiluAndMulInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/softmax/kernel.h b/src/native/cuda/mars/ops/softmax/kernel.h
new file mode 100644
index 000000000..8b7ffa2a9
--- /dev/null
+++ b/src/native/cuda/mars/ops/softmax/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SOFTMAX_KERNEL_H_
+#define INFINI_OPS_MARS_SOFTMAX_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/softmax/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Softmax, Device::Type::kMars>
+    : public CudaSoftmax<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSoftmax<Runtime<Device::Type::kMars>>::CudaSoftmax;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/softmax_infinilm/kernel.h b/src/native/cuda/mars/ops/softmax_infinilm/kernel.h
new file mode 100644
index 000000000..f402e9d3c
--- /dev/null
+++ b/src/native/cuda/mars/ops/softmax_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SOFTMAX_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_SOFTMAX_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/softmax_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<SoftmaxInfinilm, Device::Type::kMars>
+    : public CudaSoftmaxInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSoftmaxInfinilm<Runtime<Device::Type::kMars>>::CudaSoftmaxInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/swiglu/kernel.h b/src/native/cuda/mars/ops/swiglu/kernel.h
new file mode 100644
index 000000000..4155454cc
--- /dev/null
+++ b/src/native/cuda/mars/ops/swiglu/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_SWIGLU_KERNEL_H_
+#define INFINI_OPS_MARS_SWIGLU_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/swiglu/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<Swiglu, Device::Type::kMars>
+    : public CudaSwiglu<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaSwiglu<Runtime<Device::Type::kMars>>::CudaSwiglu;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/topk_softmax/kernel.h b/src/native/cuda/mars/ops/topk_softmax/kernel.h
new file mode 100644
index 000000000..a4468c735
--- /dev/null
+++ b/src/native/cuda/mars/ops/topk_softmax/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_TOPK_SOFTMAX_KERNEL_H_
+#define INFINI_OPS_MARS_TOPK_SOFTMAX_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/topk_softmax/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<TopkSoftmax, Device::Type::kMars, 0>
+    : public CudaTopkSoftmax<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaTopkSoftmax<Runtime<Device::Type::kMars>>::CudaTopkSoftmax;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h b/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h
new file mode 100644
index 000000000..687675804
--- /dev/null
+++ b/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h
@@ -0,0 +1,22 @@
+#ifndef INFINI_OPS_MARS_TOPKSOFTMAX_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_TOPKSOFTMAX_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/topksoftmax_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<TopksoftmaxInfinilm, Device::Type::kMars>
+    : public CudaTopksoftmaxInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaTopksoftmaxInfinilm<
+      Runtime<Device::Type::kMars>>::CudaTopksoftmaxInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/ops/zeros_infinilm/kernel.h b/src/native/cuda/mars/ops/zeros_infinilm/kernel.h
new file mode 100644
index 000000000..9b677ac7e
--- /dev/null
+++ b/src/native/cuda/mars/ops/zeros_infinilm/kernel.h
@@ -0,0 +1,21 @@
+#ifndef INFINI_OPS_MARS_ZEROS_INFINILM_KERNEL_H_
+#define INFINI_OPS_MARS_ZEROS_INFINILM_KERNEL_H_
+
+#include <utility>
+
+#include "native/cuda/mars/caster.cuh"
+#include "native/cuda/mars/runtime_.h"
+#include "native/cuda/ops/zeros_infinilm/kernel.h"
+
+namespace infini::ops {
+
+template <>
+class Operator<ZerosInfinilm, Device::Type::kMars>
+    : public CudaZerosInfinilm<Runtime<Device::Type::kMars>> {
+ public:
+  using CudaZerosInfinilm<Runtime<Device::Type::kMars>>::CudaZerosInfinilm;
+};
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/native/cuda/mars/runtime_.h b/src/native/cuda/mars/runtime_.h
new file mode 100644
index 000000000..ee0bbcdbd
--- /dev/null
+++ b/src/native/cuda/mars/runtime_.h
@@ -0,0 +1,9 @@
+#ifndef INFINI_OPS_MARS_RUNTIME_H_
+#define INFINI_OPS_MARS_RUNTIME_H_
+
+#include <infini/rt/mars/runtime_.h>
+
+#include "native/cuda/mars/runtime_utils.h"
+#include "runtime.h"
+
+#endif
diff --git a/src/native/cuda/mars/runtime_utils.h b/src/native/cuda/mars/runtime_utils.h
new file mode 100644
index 000000000..35c6c02bc
--- /dev/null
+++ b/src/native/cuda/mars/runtime_utils.h
@@ -0,0 +1,15 @@
+#ifndef INFINI_OPS_MARS_RUNTIME_UTILS_H_
+#define INFINI_OPS_MARS_RUNTIME_UTILS_H_
+
+#include "native/cuda/mars/device_property.h"
+#include "native/cuda/runtime_utils.h"
+
+namespace infini::ops {
+
+template <>
+struct RuntimeUtils<Device::Type::kMars>
+    : CudaRuntimeUtils<QueryMaxThreadsPerBlock> {};
+
+}  // namespace infini::ops
+
+#endif