diff --git a/CMakeLists.txt b/CMakeLists.txt
index 293b15dce..26ae3ba80 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -25,6 +25,8 @@ option(WITH_TORCH "Enable PyTorch C++ backend" OFF)
 
 option(WITH_NINETOOTHED "Enable NineToothed-generated kernels" OFF)
 
+option(WITH_TRITON "Enable Triton-generated kernels" OFF)
+
 # Custom `AscendC` kernels under `src/native/ascend/custom/`. `ON` by default
 # so CI and routine dev builds always exercise `implementation_index=1/2`
 # for `RmsNorm` / `AddRmsNorm`. Gated by `WITH_ASCEND` in
@@ -334,6 +336,10 @@ if(WITH_NINETOOTHED)
     set(NINETOOTHED_PYTHON_EXECUTABLE "" CACHE FILEPATH "Python executable used to run NineToothed code generation")
 endif()
 
+if(WITH_TRITON AND NOT WITH_NVIDIA)
+    message(FATAL_ERROR "`WITH_TRITON` temporarily requires `WITH_NVIDIA=ON` because the Triton backend temporarily targets CUDA.")
+endif()
+
 if(WITH_NVIDIA)
     add_compile_definitions(WITH_NVIDIA=1)
     enable_language(CUDA)
diff --git a/scripts/generate_wrappers.py b/scripts/generate_wrappers.py
index f69ba1396..8dc140e12 100644
--- a/scripts/generate_wrappers.py
+++ b/scripts/generate_wrappers.py
@@ -419,6 +419,8 @@ def __init__(self, name, constructors, calls):
 
         self.calls = calls
 
+        self.impl_paths = []
+
 
 def _find_optional_tensor_params(op_name):
     """Return a set of parameter names declared as `std::optional<Tensor>` in
@@ -580,6 +582,56 @@ def _is_data_type_spelling(spelling):
 
     return spelling.rsplit("::", maxsplit=1)[-1] == "DataType"
 
+def _uses_config_extension(impl_paths):
+    pattern = re.compile(r"\bJitConfig\b")
+    for path in impl_paths:
+        try:
+            if pattern.search(path.read_text()):
+                return True
+        except (OSError, UnicodeDecodeError):
+            pass
+    return False
+
+
+def _generate_triton_jit_config_parser():
+    return textwrap.dedent("""\
+    inline std::shared_ptr<Config> ConfigFromPyDict(const py::dict& config_dict) {
+      if (config_dict.contains("autotune")) {
+        auto config = std::make_shared<AutotuneConfig>();
+        py::dict autotune_dict = config_dict["autotune"].cast<py::dict>();
+        if (autotune_dict.contains("warmup")) config->warmup = autotune_dict["warmup"].cast<unsigned>();
+        if (autotune_dict.contains("rep")) config->rep = autotune_dict["rep"].cast<unsigned>();
+        if (autotune_dict.contains("key")) {
+          for (auto k : autotune_dict["key"].cast<py::list>())
+            config->key.push_back(k.cast<std::string>());
+        }
+        if (autotune_dict.contains("configs")) {
+          for (auto candidate : autotune_dict["configs"].cast<py::list>()) {
+            JitConfig candidate_config;
+            py::dict candidate_dict = candidate.cast<py::dict>();
+            if (candidate_dict.contains("num_warps")) candidate_config.num_warps = candidate_dict["num_warps"].cast<unsigned>();
+            if (candidate_dict.contains("num_stages")) candidate_config.num_stages = candidate_dict["num_stages"].cast<unsigned>();
+            for (auto item : candidate_dict) {
+              std::string key = item.first.cast<std::string>();
+              if (key != "num_warps" && key != "num_stages")
+                candidate_config.constexprs.emplace_back(key, item.second.cast<int>());
+            }
+            config->candidates.push_back(std::move(candidate_config));
+          }
+        }
+        return config;
+      }
+      auto config = std::make_shared<JitConfig>();
+      if (config_dict.contains("num_warps")) config->num_warps = config_dict["num_warps"].cast<unsigned>();
+      if (config_dict.contains("num_stages")) config->num_stages = config_dict["num_stages"].cast<unsigned>();
+      for (auto item : config_dict) {
+        std::string key = item.first.cast<std::string>();
+        if (key != "num_warps" && key != "num_stages")
+          config->constexprs.emplace_back(key, item.second.cast<int>());
+      }
+      return config;
+    }""")
+
 
 def _generate_pybind11(operator):
     optional_tensor_params = _find_optional_tensor_params(operator.name)
@@ -774,7 +826,7 @@ def _generate_py_args(node):
 
         return ", ".join(parts)
 
-    def _generate_call(op_name, call, method=True):
+    def _generate_call(op_name, call, method=True, uses_config=False):
         call_params = _generate_params(call)
         call_args = _generate_arguments(call)
 
@@ -793,12 +845,23 @@ def _generate_call(op_name, call, method=True):
                 call_args = _generate_arguments(
                     call, first_tensor_arg, converted_first_tensor_name
                 )
+            extra_params = ""
+            extra_config_init = ""
+            extra_pybind = ""
+            if uses_config:
+                extra_params = ", std::optional<py::dict> config_dict"
+                extra_config_init = (
+                    "    if (config_dict.has_value()) {\n"
+                    "      config.set_extension(ConfigFromPyDict(*config_dict));\n"
+                    "    }\n"
+                )
+                extra_pybind = ', py::arg("config") = py::none()'
+
             params = (
                 f"{call_params}, std::uintptr_t stream, "
-                "std::optional<std::size_t> implementation_index"
+                f"std::optional<std::size_t> implementation_index{extra_params}"
                 if call_params
-                else "std::uintptr_t stream, "
-                "std::optional<std::size_t> implementation_index"
+                else f"std::uintptr_t stream, std::optional<std::size_t> implementation_index{extra_params}"
             )
             py_args = _generate_py_args(call)
             py_args_str = f"{py_args}, " if py_args else ""
@@ -806,6 +869,14 @@ def _generate_call(op_name, call, method=True):
                 call, converted_first_tensor_name
             )
 
+            if uses_config:
+                dispatch = (
+                    f"    auto op = generated_dispatch::Make{symbol_name}(config, {call_args});\n"
+                    f"    (*op)(handle, {call_args});"
+                )
+            else:
+                dispatch = f"    return generated_dispatch::Call{symbol_name}(handle, config, {call_args});"
+
             return (
                 f'  m.def("{op_name}", []({params}) {{\n'
                 f"    [[maybe_unused]] HostRangeScope host_range_binding_body{{\n"
@@ -822,8 +893,9 @@ def _generate_call(op_name, call, method=True):
                 f"      config.set_implementation_index(\n"
                 f"          {default_impl_index});\n"
                 f"    }}\n"
-                f"    return generated_dispatch::Call{symbol_name}(handle, config, {call_args});\n"
-                f'  }}, {py_args_str}py::kw_only(), py::arg("stream") = 0, py::arg("implementation_index") = py::none());'
+                f"{extra_config_init}"
+                f"{dispatch}\n"
+                f'  }}, {py_args_str}py::kw_only(), py::arg("stream") = 0, py::arg("implementation_index") = py::none(){extra_pybind});'
             )
 
         # The first lambda parameter is conventionally named `self`, but
@@ -870,9 +942,21 @@ def _overload_order_key(node):
 
     inits = "\n".join(_generate_init(constructor) for constructor in constructors)
     calls = "\n".join(_generate_call(operator.name, call) for call in operator_calls)
+
+    supports_triton = _uses_config_extension(operator.impl_paths)
     callers = "\n".join(
-        _generate_call(operator.name, call, method=False) for call in operator_calls
+        _generate_call(operator.name, call, method=False, uses_config=supports_triton)
+        for call in operator_calls
     )
+    if supports_triton:
+        jit_include = (
+            '\n#include "triton/jit/jit.h"\n\n'
+            "namespace infini::ops {\n\n"
+            + _generate_triton_jit_config_parser()
+            + "\n\n}  // namespace infini::ops\n"
+        )
+    else:
+        jit_include = ""
 
     return f"""#ifndef INFINI_OPS_BINDINGS_{op_name.upper()}_H_
 #define INFINI_OPS_BINDINGS_{op_name.upper()}_H_
@@ -886,7 +970,7 @@ def _overload_order_key(node):
 #include "generated/bindings/generated_dispatch.h"
 #include "handle.h"
 #include "host_range_profiler.h"
-#include "pybind11_utils.h"
+#include "pybind11_utils.h"{jit_include}
 
 namespace py = pybind11;
 
@@ -1252,9 +1336,12 @@ def _append_optional_params(prefix, params):
 
     emitted_make_params = set()
 
-    for constructor in operator.constructors:
-        params = _generate_params(constructor)
-        args = _generate_arguments(constructor)
+    make_nodes = list(operator.constructors)
+    if _uses_config_extension(operator.impl_paths):
+        make_nodes.extend(operator.calls)
+    for node in make_nodes:
+        params = _generate_params(node)
+        args = _generate_arguments(node)
         make_params = _append_optional_params("const Config& config", params)
 
         if make_params in emitted_make_params:
@@ -1721,13 +1808,15 @@ def _filter_ops(ops, op_allowlist, *, strict=False):
     return {op_name: ops[op_name] for op_name in op_allowlist if op_name in ops}
 
 
-def _get_all_ops(devices, with_torch=False, with_ninetoothed=False):
+def _get_all_ops(devices, with_torch=False, with_ninetoothed=False, with_triton=False):
     scan_dirs = set(devices)
 
     if with_torch:
         scan_dirs.add("torch")
     if with_ninetoothed:
         scan_dirs.add("ninetoothed")
+    if with_triton:
+        scan_dirs.add("triton")
 
     ops = {}
 
@@ -1776,6 +1865,7 @@ def _generate_op_artifacts(item):
     op_name, impl_paths = item
     extractor = _OperatorExtractor()
     operator = extractor(op_name)
+    operator.impl_paths = impl_paths
     header_name = f"{op_name}.h"
     legacy_c_source, legacy_c_header = _generate_legacy_c(operator, impl_paths)
     dispatch_declarations, dispatch_definitions = _generate_generated_dispatch_entries(
@@ -1940,6 +2030,12 @@ def _dispatch_gen_batch_size():
         help="Fail if `--ops` contains operators unavailable for the active devices.",
     )
 
+    parser.add_argument(
+        "--with-triton",
+        action="store_true",
+        help="Include Triton backend implementations.",
+    )
+
     args = parser.parse_args()
 
     for directory in (_BINDINGS_DIR, _GENERATED_SRC_DIR, _INCLUDE_DIR):
@@ -1954,6 +2050,7 @@ def _dispatch_gen_batch_size():
             args.devices,
             with_torch=args.with_torch,
             with_ninetoothed=args.with_ninetoothed,
+            with_triton=args.with_triton,
         )
 
     ops = _filter_ops(
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 1edeebebd..cf2884cfa 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -189,6 +189,19 @@ if(WITH_NINETOOTHED)
     target_sources(infiniops PRIVATE ${INFINI_OPS_NINETOOTHED_SOURCES})
 endif()
 
+if(WITH_TRITON)
+    find_package(Python COMPONENTS Interpreter Development REQUIRED)
+    find_package(pybind11 CONFIG REQUIRED)
+
+    target_compile_definitions(infiniops PUBLIC WITH_TRITON=1
+        TRITON_JIT_CACHE_DIR="/tmp/triton_jit_cache")
+    target_include_directories(infiniops PRIVATE ${pybind11_INCLUDE_DIRS})
+    target_link_libraries(infiniops PRIVATE pybind11::embed Python::Python)
+    target_sources(infiniops PRIVATE triton/jit/compiler.cc)
+    file(GLOB_RECURSE TRITON_SOURCES CONFIGURE_DEPENDS "triton/ops/*/*.cc")
+    target_sources(infiniops PRIVATE ${TRITON_SOURCES})
+endif()
+
 if(WITH_ILUVATAR)
     set(ILUVATAR_PATTERNS
         "native/cuda/*.cc"
@@ -840,6 +853,10 @@ if(GENERATE_OPERATOR_CALL_INSTANTIATIONS OR GENERATE_PYTHON_BINDINGS)
         list(APPEND GENERATOR_ARGS --with-ninetoothed)
     endif()
 
+    if(WITH_TRITON)
+        list(APPEND GENERATOR_ARGS --with-triton)
+    endif()
+
     execute_process(
         COMMAND ${CMAKE_COMMAND} -E env
                 INFINI_RT_INCLUDE_DIRS=${INFINI_RT_INCLUDE_DIRS_ENV}
@@ -1190,6 +1207,7 @@ if(GENERATE_PYTHON_BINDINGS)
         target_include_directories(ops PRIVATE
             ${INFINI_OPS_NINETOOTHED_INCLUDE_DIRS})
     endif()
+
     target_link_libraries(ops PRIVATE infiniops)
 
     # Cambricon generated dispatch is compiled into the Python extension and
@@ -1244,6 +1262,19 @@ if(GENERATE_PYTHON_BINDINGS)
         install(FILES "${PROJECT_SOURCE_DIR}/generated/torch_ops_metadata.json"
             DESTINATION .)
     endif()
+
+    if(WITH_TRITON)
+        # Ship the JIT compiler and kernel sources so Triton JIT operators
+        # can compile kernels at runtime. `compile.py` uses `__file__` to
+        # locate `ops/` relative to itself; both must live under `triton/`.
+        install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/triton/jit/compile.py"
+                DESTINATION triton/jit)
+        install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/triton/ops/"
+                DESTINATION triton/ops
+                FILES_MATCHING
+                PATTERN "*.py"
+                PATTERN "build.py" EXCLUDE)
+    endif()
 endif()
 
 install(TARGETS infiniops
diff --git a/src/config.h b/src/config.h
index a8b59a4fd..15bb430ca 100644
--- a/src/config.h
+++ b/src/config.h
@@ -2,6 +2,7 @@
 #define INFINI_OPS_CONFIG_H_
 
 #include <cstddef>
+#include <memory>
 
 namespace infini::ops {
 
@@ -13,8 +14,15 @@ class Config {
     implementation_index_ = implementation_index;
   }
 
+  void set_extension(std::shared_ptr<Config> extension) {
+    extension_ = std::move(extension);
+  }
+
+  std::shared_ptr<Config> extension() const { return extension_; }
+
  private:
   std::size_t implementation_index_{0};
+  std::shared_ptr<Config> extension_{};
 };
 
 }  // namespace infini::ops
diff --git a/src/driver.h b/src/driver.h
new file mode 100644
index 000000000..9ea6a79f5
--- /dev/null
+++ b/src/driver.h
@@ -0,0 +1,16 @@
+#ifndef INFINI_OPS_DRIVER_H_
+#define INFINI_OPS_DRIVER_H_
+
+#include <infini/rt.h>
+
+namespace infini::ops {
+
+template <Device::Type device_type>
+using Driver = infini::rt::driver::Driver<device_type>;
+
+template <typename Derived>
+using DeviceDriver = infini::rt::driver::DeviceDriver<Derived>;
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/triton/jit/base.h b/src/triton/jit/base.h
new file mode 100644
index 000000000..02bdd785d
--- /dev/null
+++ b/src/triton/jit/base.h
@@ -0,0 +1,126 @@
+#ifndef INFINI_OPS_TRITON_JIT_BASE_H_
+#define INFINI_OPS_TRITON_JIT_BASE_H_
+
+#include <cassert>
+#include <string>
+#include <vector>
+
+#include "config.h"
+#include "device.h"
+#include "driver.h"
+
+namespace infini::ops {
+
+// ---- types ----
+
+struct JitConfig : Config {
+  JitConfig() = default;
+
+  JitConfig(unsigned num_warps, unsigned num_stages,
+            std::vector<std::pair<std::string, int>> constexprs)
+      : num_warps(num_warps),
+        num_stages(num_stages),
+        constexprs(std::move(constexprs)) {}
+
+  bool autotune = false;
+
+  unsigned num_warps = 4;
+
+  unsigned num_stages = 3;
+
+  std::vector<std::pair<std::string, int>> constexprs;
+
+  int At(const std::string& key) const {
+    for (const auto& [k, v] : constexprs)
+      if (k == key) return v;
+    assert(false && "`constexpr` not found");
+    return 0;
+  }
+
+  void ApplyDefaults(const JitConfig& defaults) {
+    for (const auto& [dk, dv] : defaults.constexprs) {
+      bool found = false;
+      for (const auto& [k, v] : constexprs)
+        if (k == dk) {
+          found = true;
+          break;
+        }
+      if (!found) constexprs.push_back({dk, dv});
+    }
+  }
+};
+
+struct AutotuneConfig : public JitConfig {
+  AutotuneConfig() { autotune = true; }
+
+  std::vector<std::string> key;
+
+  std::vector<JitConfig> candidates;
+
+  int warmup = 25;
+
+  int rep = 100;
+};
+
+struct Grid {
+  unsigned x = 1;
+
+  unsigned y = 1;
+
+  unsigned z = 1;
+};
+
+struct TargetInfo {
+  std::string type;
+
+  int id = 0;
+
+  int arch = 0;
+
+  int warp_size = 0;
+};
+
+struct KernelMeta {
+  std::string name;
+
+  std::string binary_ext;
+
+  unsigned shared = 0;
+
+  unsigned num_warps = 0;
+
+  int global_scratch_size = 0;
+
+  int profile_scratch_size = 0;
+};
+
+// ---- declarations ----
+
+bool CompilerInit();
+
+int CompileKernel(const TargetInfo& target, const char* op_name,
+                  const char* out_prefix, int num_warps, int num_stages,
+                  const char* signature);
+
+template <Device::Type kDev>
+int LaunchKernel(const char* op_name, const char* signature_str, void* stream,
+                 Grid grid, const JitConfig& config, void** args);
+
+template <Device::Type kDev>
+typename Driver<kDev>::Function GetKernel(const char* op_name,
+                                          const char* signature_str,
+                                          void* stream, const JitConfig& config,
+                                          unsigned* out_shared);
+
+template <Device::Type kDev>
+TargetInfo CurrentTarget();
+
+JitConfig AutotuneBench(const char* op_name,
+                        const std::vector<JitConfig>& configs,
+                        const std::string& sig, const std::vector<void*>& ptrs,
+                        const std::vector<Grid>& grids, int warmup, int rep,
+                        const char* key, const TargetInfo& target);
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/triton/jit/cache.h b/src/triton/jit/cache.h
new file mode 100644
index 000000000..3a6a68c2c
--- /dev/null
+++ b/src/triton/jit/cache.h
@@ -0,0 +1,258 @@
+#ifndef INFINI_OPS_TRITON_JIT_CACHE_H_
+#define INFINI_OPS_TRITON_JIT_CACHE_H_
+
+#include <cstdio>
+#include <cstdlib>
+#include <mutex>
+#include <sstream>
+#include <string>
+#include <unordered_map>
+
+#include "base.h"
+
+namespace infini::ops {
+
+// ---- file helpers ----
+
+inline bool FileExists(const char* path) {
+  FILE* f = fopen(path, "rb");
+  if (f != nullptr) {
+    fclose(f);
+    return true;
+  }
+  return false;
+}
+
+inline std::string ReadFile(const char* path) {
+  FILE* f = fopen(path, "rb");
+  if (f == nullptr) return {};
+  fseek(f, 0, SEEK_END);
+  long sz = ftell(f);
+  if (sz < 0) {
+    fclose(f);
+    return {};
+  }
+  fseek(f, 0, SEEK_SET);
+  std::string buf(static_cast<size_t>(sz), '\0');
+  size_t nread = fread(buf.data(), 1, static_cast<size_t>(sz), f);
+  fclose(f);
+  buf.resize(nread);
+  return buf;
+}
+
+// ---- JSON field extraction ----
+
+inline int JsonGetInt(const std::string& json, const char* key,
+                      int fallback = 0) {
+  std::string pat = std::string("\"") + key + "\":";
+  auto pos = json.find(pat);
+  if (pos == std::string::npos) return fallback;
+  pos += pat.size();
+  while (pos < json.size() && (json[pos] == ' ' || json[pos] == '\t')) pos++;
+  return std::atoi(json.c_str() + pos);
+}
+
+inline std::string JsonGetString(const std::string& json, const char* key,
+                                 const char* fallback) {
+  std::string pat = std::string("\"") + key + "\":";
+  auto pos = json.find(pat);
+  if (pos == std::string::npos) return fallback;
+  pos += pat.size();
+  while (pos < json.size() && (json[pos] == ' ' || json[pos] == '\t')) pos++;
+  if (pos >= json.size() || json[pos] != '"') return fallback;
+  pos++;
+  auto end = json.find('"', pos);
+  if (end == std::string::npos) return fallback;
+  return json.substr(pos, end - pos);
+}
+
+// ---- key generation ----
+
+inline std::string GenerateDesc(const char* op, const char* sig,
+                                unsigned num_warps, unsigned num_stages,
+                                int arch) {
+  return std::string(op) + "|" + sig + "|" + std::to_string(num_warps) + "|" +
+         std::to_string(num_stages) + "|sm" + std::to_string(arch);
+}
+
+inline std::string CacheMemKey(const char* op_name, const char* signature_str,
+                               unsigned num_warps, unsigned num_stages,
+                               int arch, int dev_id) {
+  return GenerateDesc(op_name, signature_str, num_warps, num_stages, arch) +
+         "|dev" + std::to_string(dev_id);
+}
+
+inline std::string CacheFileKey(const char* op_name, const char* signature_str,
+                                unsigned num_warps, unsigned num_stages,
+                                int arch) {
+  return std::to_string(std::hash<std::string>{}(
+      GenerateDesc(op_name, signature_str, num_warps, num_stages, arch)));
+}
+
+// ---- artifact reader ----
+
+inline bool ReadArtifacts(const std::string& out_prefix, KernelMeta* meta,
+                          std::string* binary_data) {
+  std::string meta_path = out_prefix + ".json";
+  std::string meta_json = ReadFile(meta_path.c_str());
+  if (meta_json.empty()) return false;
+
+  meta->name = JsonGetString(meta_json, "name", "");
+  meta->binary_ext = JsonGetString(meta_json, "binary_ext", "");
+  if (meta->binary_ext.empty()) meta->binary_ext = "cubin";
+  meta->shared = JsonGetInt(meta_json, "shared");
+  meta->num_warps = JsonGetInt(meta_json, "num_warps");
+  meta->global_scratch_size = JsonGetInt(meta_json, "global_scratch_size");
+  meta->profile_scratch_size = JsonGetInt(meta_json, "profile_scratch_size");
+
+  std::string binary_path = out_prefix + "." + meta->binary_ext;
+  *binary_data = ReadFile(binary_path.c_str());
+  return !binary_data->empty();
+}
+
+// ---- kernel cache ----
+
+template <Device::Type kDev>
+struct KernelCacheEntry {
+  typename Driver<kDev>::Function func;
+
+  unsigned shared;
+};
+
+template <Device::Type kDev>
+struct KernelCache {
+  std::mutex mutex;
+
+  std::unordered_map<std::string, KernelCacheEntry<kDev>> map;
+};
+
+template <Device::Type kDev>
+KernelCache<kDev>& GetKernelCache() {
+  static KernelCache<kDev> c;
+  return c;
+}
+
+template <Device::Type kDev>
+bool KernelCacheLookup(const std::string& key, KernelCacheEntry<kDev>* out) {
+  auto& c = GetKernelCache<kDev>();
+  std::lock_guard<std::mutex> lk(c.mutex);
+  auto it = c.map.find(key);
+  if (it == c.map.end()) return false;
+  *out = it->second;
+  return true;
+}
+
+template <Device::Type kDev>
+void KernelCacheInsert(const std::string& key, KernelCacheEntry<kDev> entry) {
+  auto& c = GetKernelCache<kDev>();
+  std::lock_guard<std::mutex> lk(c.mutex);
+  c.map[key] = entry;
+}
+
+template <Device::Type kDev>
+struct CacheQueryResult {
+  bool mem_hit;
+
+  typename Driver<kDev>::Function func;
+
+  unsigned shared;
+
+  std::string out_prefix;
+
+  std::string mem_key;
+};
+
+template <Device::Type kDev>
+CacheQueryResult<kDev> CacheQuery(const char* op, const char* sig,
+                                  unsigned num_warps, unsigned num_stages,
+                                  int arch, int dev_id) {
+  auto mem_key = CacheMemKey(op, sig, num_warps, num_stages, arch, dev_id);
+  KernelCacheEntry<kDev> entry;
+  if (KernelCacheLookup<kDev>(mem_key, &entry))
+    return {true, entry.func, entry.shared, "", mem_key};
+  auto desc = GenerateDesc(op, sig, num_warps, num_stages, arch);
+  return {false, nullptr, 0,
+          std::string(TRITON_JIT_CACHE_DIR) + "/" +
+              std::to_string(std::hash<std::string>{}(desc)),
+          mem_key};
+}
+
+// ---- autotune cache ----
+
+struct AutotuneCache {
+  std::mutex mutex;
+
+  std::unordered_map<std::string, JitConfig> map;
+};
+
+inline AutotuneCache& GetAutotuneCache() {
+  static AutotuneCache c;
+  return c;
+}
+
+inline std::string AutotuneCacheFilePath(const std::string& key) {
+  return std::string{TRITON_JIT_CACHE_DIR} + "/" +
+         std::to_string(std::hash<std::string>{}(key)) + ".autotune";
+}
+
+inline std::string SerializeConfig(const JitConfig& config) {
+  std::string s = std::to_string(config.num_warps) + " " +
+                  std::to_string(config.num_stages);
+  for (const auto& [name, val] : config.constexprs)
+    s += "\n" + name + " " + std::to_string(val);
+  return s;
+}
+
+inline bool DeserializeConfig(const std::string& content, JitConfig* out) {
+  std::istringstream iss(content);
+  std::string line;
+  if (!std::getline(iss, line)) return false;
+  std::istringstream head(line);
+  if (!(head >> out->num_warps >> out->num_stages)) return false;
+  out->constexprs.clear();
+  while (std::getline(iss, line)) {
+    std::istringstream ls(line);
+    std::string name;
+    int val;
+    if (ls >> name >> val) out->constexprs.push_back({name, val});
+  }
+  return true;
+}
+
+inline bool AutotuneCacheLookup(const std::string& key, JitConfig* out) {
+  auto& c = GetAutotuneCache();
+  std::lock_guard<std::mutex> lk(c.mutex);
+  auto it = c.map.find(key);
+  if (it != c.map.end()) {
+    *out = it->second;
+    return true;
+  }
+  std::string path = AutotuneCacheFilePath(key);
+  if (FileExists(path.c_str())) {
+    JitConfig parsed;
+    if (DeserializeConfig(ReadFile(path.c_str()), &parsed)) {
+      c.map[key] = parsed;
+      *out = parsed;
+      return true;
+    }
+  }
+  return false;
+}
+
+inline void AutotuneCacheInsert(const std::string& key,
+                                const JitConfig& config) {
+  auto& c = GetAutotuneCache();
+  std::lock_guard<std::mutex> lk(c.mutex);
+  c.map[key] = config;
+  std::string path = AutotuneCacheFilePath(key);
+  std::string content = SerializeConfig(config);
+  FILE* f = fopen(path.c_str(), "w");
+  if (f) {
+    fwrite(content.data(), 1, content.size(), f);
+    fclose(f);
+  }
+}
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/triton/jit/compile.py b/src/triton/jit/compile.py
new file mode 100644
index 000000000..c1a6d8f77
--- /dev/null
+++ b/src/triton/jit/compile.py
@@ -0,0 +1,163 @@
+import importlib.util
+import json
+from pathlib import Path
+
+import triton
+import triton.backends
+
+_JIT_DIR = Path(__file__).resolve().parent
+_OPS_DIR = _JIT_DIR.parent / "ops"
+
+_TRITON_BACKEND = {
+    "nvidia": "cuda",
+}
+
+
+def _make_target(device, device_id, arch, warp_size):
+    triton.runtime.driver.set_active(triton.backends.backends[device].driver())
+    triton.runtime.driver.active.set_current_device(device_id)
+    return triton.backends.compiler.GPUTarget(_TRITON_BACKEND[device], arch, warp_size)
+
+
+def _do_compile(
+    op_name,
+    out_prefix,
+    num_warps,
+    num_stages,
+    device_id,
+    signature,
+    device,
+    arch,
+    warp_size,
+):
+
+    source_path = _OPS_DIR / f"{op_name}/{op_name}.py"
+    spec = importlib.util.spec_from_file_location(source_path.stem, source_path)
+    mod = importlib.util.module_from_spec(spec)
+    spec.loader.exec_module(mod)
+    fn = getattr(mod, "kernel")
+    while not isinstance(fn, triton.runtime.JITFunction):
+        fn = fn.fn
+
+    sig_parts = [p.strip() for p in signature.split(",")] if signature else []
+    assert len(sig_parts) == len(fn.arg_names), (
+        f"signature length {len(sig_parts)} != kernel param count {len(fn.arg_names)}"
+    )
+
+    sig_dict = {}
+    const_dict = {}
+    attr_dict = {}
+
+    constexprs = {}
+    for part in sig_parts:
+        if "=" in part:
+            name, val = part.split("=", 1)
+            constexprs[name.strip()] = int(val)
+
+    for i, (name, param, part) in enumerate(zip(fn.arg_names, fn.params, sig_parts)):
+        if param.is_constexpr:
+            const_dict[(i,)] = constexprs[name]
+            sig_dict[name] = "constexpr"
+        elif part.endswith(":1"):
+            const_dict[(i,)] = 1
+            sig_dict[name] = "constexpr"
+        elif part.endswith(":16"):
+            sig_dict[name] = part[:-3]
+            attr_dict[(i,)] = [["tt.divisibility", 16]]
+        else:
+            sig_dict[name] = part
+
+    src = triton.compiler.ASTSource(
+        fn=fn, signature=sig_dict, constexprs=const_dict, attrs=attr_dict
+    )
+
+    target = _make_target(device, device_id, arch, warp_size)
+    ccinfo = triton.compile(
+        src,
+        target=target,
+        options={"num_warps": num_warps, "num_stages": num_stages},
+    )
+
+    Path(out_prefix).parent.mkdir(parents=True, exist_ok=True)
+    backend = triton.compiler.make_backend(target)
+    bin_ext = backend.binary_ext
+    binary = ccinfo.asm[bin_ext]
+    with open(out_prefix + "." + bin_ext, "wb") as f:
+        f.write(binary)
+
+    meta = {
+        "name": getattr(ccinfo.metadata, "name", fn.__name__),
+        "binary_ext": bin_ext,
+        "shared": getattr(ccinfo.metadata, "shared", 0),
+        "num_warps": getattr(ccinfo.metadata, "num_warps", num_warps),
+        "device": device,
+        "target_backend": target.backend,
+        "arch": str(target.arch),
+        "global_scratch_size": getattr(ccinfo.metadata, "global_scratch_size", 0),
+        "profile_scratch_size": getattr(ccinfo.metadata, "profile_scratch_size", 0),
+        "op_name": op_name,
+        "signature": signature,
+    }
+    with open(out_prefix + ".json", "w") as f:
+        json.dump(meta, f)
+
+
+def _load_kernel_fn(op_name):
+    source_path = _OPS_DIR / f"{op_name}/{op_name}.py"
+    spec = importlib.util.spec_from_file_location(source_path.stem, source_path)
+    mod = importlib.util.module_from_spec(spec)
+    spec.loader.exec_module(mod)
+    fn = getattr(mod, "kernel")
+    while not isinstance(fn, triton.runtime.JITFunction):
+        fn = fn.fn
+    return fn
+
+
+def _do_autotune(
+    op_name,
+    configs,
+    args,
+    grids,
+    warmup,
+    rep,
+    device_id,
+    device,
+    arch,
+    warp_size,
+):
+    fn = _load_kernel_fn(op_name)
+    best_idx = 0
+    best_time = float("inf")
+    _make_target(device, device_id, arch, warp_size)
+    for i, cand in enumerate(configs):
+        constexprs = {kv[0]: kv[1] for kv in cand["constexprs"]}
+        num_warps = cand["num_warps"]
+        num_stages = cand["num_stages"]
+        grid = tuple(grids[i])
+
+        out_prefix = cand["out_prefix"]
+        _do_compile(
+            op_name,
+            out_prefix,
+            num_warps,
+            num_stages,
+            device_id,
+            cand["full_sig"],
+            device,
+            arch,
+            warp_size,
+        )
+
+        def _kernel_call(g=grid, a=args, ce=constexprs, nw=num_warps, ns=num_stages):
+            fn[g](*a, **ce, num_warps=nw, num_stages=ns)
+
+        try:
+            t = triton.testing.do_bench(
+                _kernel_call, warmup=warmup, rep=rep, quantiles=(0.5, 0.2, 0.8)
+            )[0]
+            if t < best_time:
+                best_time = t
+                best_idx = i
+        except Exception:
+            pass
+    return best_idx
diff --git a/src/triton/jit/compiler.cc b/src/triton/jit/compiler.cc
new file mode 100644
index 000000000..f53d99015
--- /dev/null
+++ b/src/triton/jit/compiler.cc
@@ -0,0 +1,151 @@
+#include <pybind11/embed.h>
+
+#include <cstdio>
+#include <mutex>
+
+#include "jit.h"
+
+namespace infini::ops {
+
+bool CompilerInit() {
+  static std::once_flag flag;
+  static bool ready = false;
+
+  std::call_once(flag, [] {
+    namespace py = pybind11;
+
+    auto setup = [] { py::module_::import("infini.triton.jit.compile"); };
+
+    if (Py_IsInitialized()) {
+      py::gil_scoped_acquire gil;
+      try {
+        setup();
+        ready = true;
+      } catch (const py::error_already_set& e) {
+        fprintf(stderr, "jit init: %s\n", e.what());
+      }
+    } else {
+      py::initialize_interpreter(false);
+      try {
+        setup();
+        ready = true;
+      } catch (const py::error_already_set& e) {
+        fprintf(stderr, "jit init: %s\n", e.what());
+      }
+      (void)PyEval_SaveThread();
+    }
+  });
+
+  return ready;
+}
+
+int CompileKernel(const TargetInfo& target, const char* op_name,
+                  const char* out_prefix, int num_warps, int num_stages,
+                  const char* signature) {
+  if (!CompilerInit()) return -1;
+
+  namespace py = pybind11;
+  py::gil_scoped_acquire gil;
+  try {
+    py::module_ mod = py::module_::import("infini.triton.jit.compile");
+    mod.attr("_do_compile")(op_name, out_prefix, num_warps, num_stages,
+                            target.id, signature, target.type, target.arch,
+                            target.warp_size);
+    return 0;
+  } catch (const py::error_already_set& e) {
+    fprintf(stderr, "jit compile: %s\n", e.what());
+    return -2;
+  }
+}
+
+JitConfig AutotuneBench(const char* op_name,
+                        const std::vector<JitConfig>& configs,
+                        const std::string& sig, const std::vector<void*>& ptrs,
+                        const std::vector<Grid>& grids, int warmup, int rep,
+                        const char* key, const TargetInfo& target) {
+  JitConfig cached;
+  if (AutotuneCacheLookup(key, &cached)) return cached;
+
+  namespace py = pybind11;
+  if (!CompilerInit()) return configs.empty() ? JitConfig{} : configs[0];
+  py::gil_scoped_acquire gil;
+  try {
+    py::module_ mod = py::module_::import("infini.triton.jit.compile");
+
+    py::list cands;
+    for (const auto& c : configs) {
+      py::dict cd;
+      cd["num_warps"] = c.num_warps;
+      cd["num_stages"] = c.num_stages;
+      py::list ce;
+      for (const auto& [k, v] : c.constexprs) {
+        py::tuple kv(2);
+        kv[0] = k;
+        kv[1] = v;
+        ce.append(kv);
+      }
+      cd["constexprs"] = ce;
+
+      std::string full_sig = sig;
+      for (const auto& [k, v] : c.constexprs)
+        full_sig += k + "=" + std::to_string(v) + ",";
+      if (!full_sig.empty() && full_sig.back() == ',') full_sig.pop_back();
+      cd["full_sig"] = full_sig;
+      cd["out_prefix"] = std::string(TRITON_JIT_CACHE_DIR) + "/" +
+                         CacheFileKey(op_name, full_sig.c_str(), c.num_warps,
+                                      c.num_stages, target.arch);
+
+      cands.append(cd);
+    }
+
+    py::list args;
+    size_t ptr_idx = 0;
+    size_t pos = 0;
+    while (pos < sig.size()) {
+      size_t comma = sig.find(',', pos);
+      std::string part = sig.substr(pos, comma - pos);
+      pos = (comma == std::string::npos) ? sig.size() : comma + 1;
+      if (part.empty()) continue;
+
+      if (part[0] == '*') {
+        uint64_t val = *static_cast<uint64_t*>(ptrs[ptr_idx++]);
+        args.append(static_cast<int64_t>(val));
+      } else if (part.find(":1") != std::string::npos) {
+        args.append(1);
+      } else {
+        uint64_t val = *static_cast<uint64_t*>(ptrs[ptr_idx++]);
+        if (part.compare(0, 4, "fp32") == 0 || part.compare(0, 3, "f32") == 0) {
+          args.append(*reinterpret_cast<float*>(&val));
+        } else if (part.compare(0, 4, "fp64") == 0) {
+          args.append(*reinterpret_cast<double*>(&val));
+        } else {
+          args.append(static_cast<int64_t>(val));
+        }
+      }
+    }
+
+    py::list grids_list;
+    for (const auto& g : grids) {
+      py::tuple t(3);
+      t[0] = g.x;
+      t[1] = g.y;
+      t[2] = g.z;
+      grids_list.append(t);
+    }
+
+    int best_idx = mod.attr("_do_autotune")(op_name, cands, args, grids_list,
+                                            warmup, rep, target.id, target.type,
+                                            target.arch, target.warp_size)
+                       .cast<int>();
+    if (best_idx < 0 || best_idx >= static_cast<int>(configs.size()))
+      best_idx = 0;
+    JitConfig winner = configs[best_idx];
+    AutotuneCacheInsert(key, winner);
+    return winner;
+  } catch (const py::error_already_set& e) {
+    fprintf(stderr, "jit autotune: %s\n", e.what());
+    return configs.empty() ? JitConfig{} : configs[0];
+  }
+}
+
+}  // namespace infini::ops
diff --git a/src/triton/jit/jit.h b/src/triton/jit/jit.h
new file mode 100644
index 000000000..ab955e881
--- /dev/null
+++ b/src/triton/jit/jit.h
@@ -0,0 +1,326 @@
+#ifndef INFINI_OPS_TRITON_JIT_H_
+#define INFINI_OPS_TRITON_JIT_H_
+
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <deque>
+#include <string>
+#include <type_traits>
+#include <vector>
+
+#include "cache.h"
+#include "data_type.h"
+#include "runtime.h"
+#include "tensor.h"
+
+namespace infini::ops {
+
+// ---- device support ----
+
+template <Device::Type kDev>
+inline constexpr bool kJitSupported = false;
+
+template <>
+inline constexpr bool kJitSupported<Device::Type::kNvidia> = true;
+
+template <typename Op, Device::Type kDev, bool = kJitSupported<kDev>>
+struct JitOperatorBase : Op {
+  using Op::Op;
+};
+
+template <typename Op, Device::Type kDev>
+struct JitOperatorBase<Op, kDev, false> {};
+
+// ---- compilation & launch ----
+
+template <Device::Type kDev>
+TargetInfo CurrentTarget() {
+  TargetInfo target;
+  target.type = Device::StringFromType(kDev);
+  int dev_id = 0;
+  if (Runtime<kDev>::GetDevice(&dev_id) != Runtime<kDev>::kSuccess)
+    return target;
+  target.id = dev_id;
+  int major = 0, minor = 0;
+  Runtime<kDev>::DeviceGetAttribute(
+      &major, Runtime<kDev>::kDevAttrComputeCapabilityMajor, dev_id);
+  Runtime<kDev>::DeviceGetAttribute(
+      &minor, Runtime<kDev>::kDevAttrComputeCapabilityMinor, dev_id);
+  target.arch = major * 10 + minor;
+  Runtime<kDev>::DeviceGetAttribute(&target.warp_size,
+                                    Runtime<kDev>::kDevAttrWarpSize, dev_id);
+  return target;
+}
+
+template <Device::Type kDev>
+bool Load(const TargetInfo& target, const char* binary_data, size_t binary_size,
+          const KernelMeta& meta, typename Driver<kDev>::Function& func,
+          typename Driver<kDev>::Module& mod) {
+  (void)binary_size;
+
+  if (Driver<kDev>::ModuleLoadData(&mod, binary_data) != Driver<kDev>::kSuccess)
+    return false;
+
+  if (Driver<kDev>::ModuleGetFunction(&func, mod, meta.name.c_str()) !=
+      Driver<kDev>::kSuccess) {
+    Driver<kDev>::ModuleUnload(mod);
+    return false;
+  }
+
+  if (meta.shared > 49152) {
+    int optin = 0;
+    Runtime<kDev>::DeviceGetAttribute(
+        &optin, Runtime<kDev>::kDevAttrMaxSharedMemoryPerBlockOptin, target.id);
+    int st = 0;
+    Driver<kDev>::FuncGetAttribute(
+        &st, Driver<kDev>::kFuncAttributeSharedSizeBytes, func);
+    if (optin < st || meta.shared > static_cast<unsigned>(optin - st)) {
+      Driver<kDev>::ModuleUnload(mod);
+      return false;
+    }
+    Driver<kDev>::FuncSetCacheConfig(func,
+                                     Driver<kDev>::kFuncCachePreferShared);
+    if (Driver<kDev>::FuncSetAttribute(
+            func, Driver<kDev>::kFuncAttributeMaxDynamicSharedSizeBytes,
+            optin - st) != Driver<kDev>::kSuccess) {
+      Driver<kDev>::ModuleUnload(mod);
+      return false;
+    }
+  }
+
+  return true;
+}
+
+template <Device::Type kDev>
+typename Driver<kDev>::Function GetKernel(const char* op_name,
+                                          const char* signature_str,
+                                          void* stream, const JitConfig& opts,
+                                          unsigned* out_shared) {
+  TargetInfo target = CurrentTarget<kDev>();
+
+  auto r = CacheQuery<kDev>(op_name, signature_str, opts.num_warps,
+                            opts.num_stages, target.arch, target.id);
+  if (r.mem_hit) {
+    *out_shared = r.shared;
+    return r.func;
+  }
+
+  KernelMeta meta;
+  std::string binary_data;
+  if (!ReadArtifacts(r.out_prefix, &meta, &binary_data)) {
+    int ret = CompileKernel(target, op_name, r.out_prefix.c_str(),
+                            opts.num_warps, opts.num_stages, signature_str);
+    if (ret != 0) return nullptr;
+    if (!ReadArtifacts(r.out_prefix, &meta, &binary_data)) return nullptr;
+  }
+
+  if (meta.global_scratch_size > 0 || meta.profile_scratch_size > 0) {
+    fprintf(stderr, "triton jit: scratch not supported yet\n");
+    return nullptr;
+  }
+
+  typename Driver<kDev>::Function func;
+  typename Driver<kDev>::Module mod;
+  if (!Load<kDev>(target, binary_data.data(), binary_data.size(), meta, func,
+                  mod))
+    return nullptr;
+
+  unsigned shared = meta.shared;
+  KernelCacheEntry<kDev> mine{func, shared};
+  KernelCacheEntry<kDev> winner;
+  if (KernelCacheLookup<kDev>(r.mem_key, &winner)) {
+    Driver<kDev>::ModuleUnload(mod);
+    func = winner.func;
+    shared = winner.shared;
+  } else {
+    KernelCacheInsert<kDev>(r.mem_key, mine);
+  }
+
+  *out_shared = shared;
+  return func;
+}
+
+template <Device::Type kDev>
+int LaunchKernel(const char* op_name, const char* signature_str, void* stream,
+                 Grid grid, const JitConfig& config, void** args) {
+  unsigned shared = 0;
+  auto func = GetKernel<kDev>(op_name, signature_str, stream, config, &shared);
+  if (!func) return -1;
+
+  TargetInfo target = CurrentTarget<kDev>();
+  return Driver<kDev>::LaunchKernel(
+      func, grid.x, grid.y, grid.z, config.num_warps * target.warp_size, 1, 1,
+      shared, static_cast<typename Driver<kDev>::Stream>(stream), args,
+      nullptr);
+}
+
+// ---- specialization helpers ----
+
+inline const char* SpecPtr(uintptr_t v) { return v % 16 == 0 ? ":16" : ""; }
+
+template <typename T>
+const char* SpecInt(T v) {
+  if (v == 1) return ":1";
+  if ((v & 15) == 0) return ":16";
+  return "";
+}
+
+// ---- `DataType` → Triton string ----
+
+inline const char* DataTypeToTritonType(DataType dt) {
+  switch (dt) {
+    case DataType::kFloat16:
+      return "fp16";
+    case DataType::kBFloat16:
+      return "bf16";
+    case DataType::kFloat32:
+      return "fp32";
+    case DataType::kFloat64:
+      return "fp64";
+    case DataType::kInt8:
+      return "i8";
+    case DataType::kInt16:
+      return "i16";
+    case DataType::kInt32:
+      return "i32";
+    case DataType::kInt64:
+      return "i64";
+    case DataType::kUInt8:
+      return "u8";
+    case DataType::kUInt16:
+      return "u16";
+    case DataType::kUInt32:
+      return "u32";
+    case DataType::kUInt64:
+      return "u64";
+  }
+  return "fp32";
+}
+
+// ---- C++ scalar type → Triton string ----
+
+template <typename T>
+const char* ScalarTypeToTritonType() {
+  if constexpr (std::is_same_v<T, float>)
+    return "fp64";
+  else if constexpr (std::is_same_v<T, double>)
+    return "fp64";
+  else if constexpr (std::is_same_v<T, bool>)
+    return "i32";
+  else if constexpr (std::is_integral_v<T>) {
+    if constexpr (sizeof(T) == 1) return std::is_signed_v<T> ? "i8" : "u8";
+    if constexpr (sizeof(T) == 2) return std::is_signed_v<T> ? "i16" : "u16";
+    if constexpr (sizeof(T) == 4) return std::is_signed_v<T> ? "i32" : "u32";
+    if constexpr (sizeof(T) == 8) return std::is_signed_v<T> ? "i64" : "u64";
+  }
+  return "i32";
+}
+
+// ---- arguments parser ----
+
+struct ArgPack {
+  std::vector<void*> ptrs;
+
+  std::deque<uint64_t> storage;
+
+  std::string sig;
+
+  template <typename T>
+  void* Store(T v) {
+    static_assert(sizeof(T) <= sizeof(uint64_t),
+                  "scalar arg wider than `uint64_t`");
+    uint64_t slot = 0;
+    std::memcpy(&slot, &v, sizeof(T));
+    storage.push_back(slot);
+    return &storage.back();
+  }
+};
+
+inline void PushArg(const Tensor& t, ArgPack& pack) {
+  auto ptr = reinterpret_cast<uintptr_t>(t.data());
+  pack.sig +=
+      std::string("*") + DataTypeToTritonType(t.dtype()) + SpecPtr(ptr) + ",";
+  pack.ptrs.push_back(pack.Store(ptr));
+}
+
+template <typename T, std::enable_if_t<std::is_integral_v<T>, int> = 0>
+void PushArg(T v, ArgPack& pack) {
+  const char* s = SpecInt(v);
+  pack.sig += std::string(ScalarTypeToTritonType<T>()) + s + ",";
+  if (std::strcmp(s, ":1") != 0) pack.ptrs.push_back(pack.Store(v));
+}
+
+inline void PushArg(float v, ArgPack& pack) {
+  pack.ptrs.push_back(pack.Store(v));
+  pack.sig += "fp32,";
+}
+
+inline void PushArg(double v, ArgPack& pack) {
+  pack.ptrs.push_back(pack.Store(v));
+  pack.sig += "fp64,";
+}
+
+// ---- launch wrapper ----
+
+template <Device::Type kDev, typename... Args>
+int LaunchJit(const char* op, void* stream, Grid grid, const JitConfig& config,
+              Args&&... args) {
+  ArgPack pack;
+  pack.sig.reserve(256);
+  (PushArg(std::forward<Args>(args), pack), ...);
+  for (const auto& [name, val] : config.constexprs)
+    pack.sig += name + "=" + std::to_string(val) + ",";
+  if (!pack.sig.empty()) pack.sig.pop_back();
+
+  void* scratch = pack.Store<uint64_t>(0);
+  pack.ptrs.push_back(scratch);
+  pack.ptrs.push_back(scratch);
+
+  return LaunchKernel<kDev>(op, pack.sig.c_str(), stream, grid, config,
+                            pack.ptrs.data());
+}
+
+template <Device::Type kDev, typename GridFn, typename... Args>
+int LaunchJitAutotune(const char* op, void* stream,
+                      const AutotuneConfig& config,
+                      const std::vector<Tensor::Size>& key,
+                      const std::vector<DataType>& dtype, GridFn grid_fn,
+                      Args&&... args) {
+  TargetInfo target = CurrentTarget<kDev>();
+
+  std::string cache_key = op;
+  for (auto d : key) cache_key += "|" + std::to_string(d);
+  for (auto dt : dtype)
+    cache_key += "|" + std::string(DataTypeToTritonType(dt));
+  cache_key += "|sm" + std::to_string(target.arch);
+
+  ArgPack pack;
+  pack.sig.reserve(256);
+  (PushArg(std::forward<Args>(args), pack), ...);
+
+  std::vector<Grid> grids;
+  grids.reserve(config.candidates.size());
+  for (const auto& c : config.candidates) grids.push_back(grid_fn(c));
+
+  JitConfig best =
+      AutotuneBench(op, config.candidates, pack.sig, pack.ptrs, grids,
+                    config.warmup, config.rep, cache_key.c_str(), target);
+
+  Grid grid = grid_fn(best);
+
+  for (const auto& [name, val] : best.constexprs)
+    pack.sig += name + "=" + std::to_string(val) + ",";
+  if (!pack.sig.empty()) pack.sig.pop_back();
+
+  void* scratch = pack.Store<uint64_t>(0);
+  pack.ptrs.push_back(scratch);
+  pack.ptrs.push_back(scratch);
+
+  return LaunchKernel<kDev>(op, pack.sig.c_str(), stream, grid, best,
+                            pack.ptrs.data());
+}
+
+}  // namespace infini::ops
+
+#endif
diff --git a/src/triton/ops/add/add.py b/src/triton/ops/add/add.py
new file mode 100644
index 000000000..f813cda47
--- /dev/null
+++ b/src/triton/ops/add/add.py
@@ -0,0 +1,53 @@
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def kernel(
+    x_ptr,
+    y_ptr,
+    out_ptr,
+    out_shape_ptr,
+    x_stride_ptr,
+    y_stride_ptr,
+    out_stride_ptr,
+    x_contig,
+    y_contig,
+    out_contig,
+    ndim,
+    n_elements,
+    alpha,
+    BLOCK_SIZE: tl.constexpr,
+):
+    pid = tl.program_id(0)
+    offsets = (pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)).to(tl.int64)
+    mask = offsets < n_elements
+
+    if (x_contig != 0) and (y_contig != 0) and (out_contig != 0):
+        x = tl.load(x_ptr + offsets, mask=mask)
+        y = tl.load(y_ptr + offsets, mask=mask)
+        tl.store(out_ptr + offsets, x + y * alpha, mask=mask)
+    else:
+        x_offs = tl.zeros([BLOCK_SIZE], dtype=tl.int64)
+        y_offs = tl.zeros([BLOCK_SIZE], dtype=tl.int64)
+        out_offs = tl.zeros([BLOCK_SIZE], dtype=tl.int64)
+        tmp = offsets
+
+        for i in range(ndim):
+            s = tl.load(out_shape_ptr + (ndim - 1 - i))
+            d = tmp % s
+            tmp = tmp // s
+            x_offs += d * tl.load(x_stride_ptr + (ndim - 1 - i))
+            y_offs += d * tl.load(y_stride_ptr + (ndim - 1 - i))
+            out_offs += d * tl.load(out_stride_ptr + (ndim - 1 - i))
+
+        if x_contig != 0:
+            x_offs = offsets
+        if y_contig != 0:
+            y_offs = offsets
+        if out_contig != 0:
+            out_offs = offsets
+
+        x = tl.load(x_ptr + x_offs, mask=mask)
+        y = tl.load(y_ptr + y_offs, mask=mask)
+        tl.store(out_ptr + out_offs, x + y * alpha, mask=mask)
diff --git a/src/triton/ops/add/jit.cc b/src/triton/ops/add/jit.cc
new file mode 100644
index 000000000..00c6b0842
--- /dev/null
+++ b/src/triton/ops/add/jit.cc
@@ -0,0 +1,96 @@
+#include "triton/ops/add/jit.h"
+
+#include <cassert>
+#include <memory>
+#include <unordered_map>
+
+#include "runtime.h"
+#include "triton/jit/jit.h"
+
+namespace infini::ops {
+
+template <Device::Type kDev>
+void Operator<Add, kDev, 10>::operator()(const Tensor input, const Tensor other,
+                                         const double alpha, Tensor out) const {
+  const int ndim = static_cast<int>(this->ndim_);
+
+  std::vector<int64_t> h_meta(4 * std::max(ndim, 1), 0);
+  for (int i = 0; i < ndim; ++i) {
+    h_meta[0 * ndim + i] = static_cast<int64_t>(this->out_shape_[i]);
+    h_meta[1 * ndim + i] = static_cast<int64_t>(this->input_strides_[i]);
+    h_meta[2 * ndim + i] = static_cast<int64_t>(this->other_strides_[i]);
+    h_meta[3 * ndim + i] = static_cast<int64_t>(this->out_strides_[i]);
+  }
+  const size_t meta_bytes = h_meta.size() * sizeof(int64_t);
+  void* d_meta = nullptr;
+  Runtime<kDev>::Malloc(&d_meta, meta_bytes);
+  Runtime<kDev>::Memcpy(d_meta, h_meta.data(), meta_bytes,
+                        Runtime<kDev>::kMemcpyHostToDevice);
+  const size_t stride_bytes = ndim * sizeof(int64_t);
+
+  std::vector<Tensor::Size> meta_shape{
+      static_cast<Tensor::Size>(std::max(ndim, 1))};
+  char* base = static_cast<char*>(d_meta);
+  Tensor d_out_shape{base + stride_bytes * 0, meta_shape, DataType::kInt64,
+                     out.device()};
+  Tensor d_input_strides{base + stride_bytes * 1, meta_shape, DataType::kInt64,
+                         out.device()};
+  Tensor d_other_strides{base + stride_bytes * 2, meta_shape, DataType::kInt64,
+                         out.device()};
+  Tensor d_out_strides{base + stride_bytes * 3, meta_shape, DataType::kInt64,
+                       out.device()};
+
+  const size_t n_elements = out.numel();
+
+  static const JitConfig defaults = DefaultConfig();
+  std::shared_ptr<Config> extension = this->config_.extension();
+  auto cfg = std::static_pointer_cast<JitConfig>(extension);
+
+  const std::unordered_map<std::string, Tensor::Size> args{
+      {"n_elements", n_elements},
+      {"ndim", ndim},
+  };
+
+  int result;
+  if (cfg && cfg->autotune) {
+    auto tune = std::static_pointer_cast<AutotuneConfig>(extension);
+    if (tune->candidates.empty()) tune->candidates = AutotuneConfigs();
+    for (auto& c : tune->candidates) c.ApplyDefaults(defaults);
+
+    auto key_names = tune->key.empty() ? DefaultKey() : tune->key;
+    std::vector<Tensor::Size> key_vals;
+    for (const auto& name : key_names) key_vals.push_back(args.at(name));
+
+    result = LaunchJitAutotune<kDev>(
+        "add", this->stream_, *tune, key_vals,
+        {input.dtype(), other.dtype(), out.dtype()},
+        [&](const JitConfig& c) {
+          int block_size = c.At("BLOCK_SIZE");
+          return Grid{static_cast<unsigned>((n_elements + block_size - 1) /
+                                            block_size)};
+        },
+        input, other, out, d_out_shape, d_input_strides, d_other_strides,
+        d_out_strides, this->is_input_contiguous_, this->is_other_contiguous_,
+        this->is_out_contiguous_, ndim, n_elements, alpha);
+  } else {
+    JitConfig config = cfg ? *cfg : defaults;
+    if (cfg) config.ApplyDefaults(defaults);
+    const int block_size = config.At("BLOCK_SIZE");
+    Grid grid{
+        static_cast<unsigned>((n_elements + block_size - 1) / block_size)};
+    result = LaunchJit<kDev>("add", this->stream_, grid, config, input, other,
+                             out, d_out_shape, d_input_strides, d_other_strides,
+                             d_out_strides, this->is_input_contiguous_,
+                             this->is_other_contiguous_,
+                             this->is_out_contiguous_, ndim, n_elements, alpha);
+  }
+
+  Runtime<kDev>::FreeAsync(
+      d_meta, static_cast<typename Runtime<kDev>::Stream>(this->stream_));
+
+  assert(result == 0 && "Triton JIT `Add` launch failed");
+}
+
+template class Operator<Add, Device::Type::kNvidia, 10>;
+
+}  // namespace infini::ops
diff --git a/src/triton/ops/add/jit.h b/src/triton/ops/add/jit.h
new file mode 100644
index 000000000..be5499072
--- /dev/null
+++ b/src/triton/ops/add/jit.h
@@ -0,0 +1,36 @@
+#ifndef INFINI_OPS_TRITON_OPS_ADD_JIT_H_
+#define INFINI_OPS_TRITON_OPS_ADD_JIT_H_
+
+#include <string>
+#include <vector>
+
+#include "base/add.h"
+#include "triton/jit/jit.h"
+
+namespace infini::ops {
+
+template <Device::Type kDev>
+class Operator<Add, kDev, 10> : public JitOperatorBase<Add, kDev> {
+ public:
+  using JitOperatorBase<Add, kDev>::JitOperatorBase;
+
+  void operator()(const Tensor input, const Tensor other, const double alpha,
+                  Tensor out) const;
+
+  static JitConfig DefaultConfig() { return {4u, 3u, {{"BLOCK_SIZE", 1024}}}; }
+
+  static std::vector<std::string> DefaultKey() { return {"n_elements"}; }
+
+  static std::vector<JitConfig> AutotuneConfigs() {
+    return {
+        {4u, 3u, {{"BLOCK_SIZE", 256}}},
+        {4u, 3u, {{"BLOCK_SIZE", 512}}},
+        {8u, 4u, {{"BLOCK_SIZE", 1024}}},
+        {8u, 4u, {{"BLOCK_SIZE", 2048}}},
+    };
+  }
+};
+
+}  // namespace infini::ops
+
+#endif