diff --git a/docsrc/py_api/kernels.rst b/docsrc/py_api/kernels.rst
index a6eda8ff4b..88f152721b 100644
--- a/docsrc/py_api/kernels.rst
+++ b/docsrc/py_api/kernels.rst
@@ -17,19 +17,22 @@ torch_tensorrt.kernels
 Overview
 --------
 
-The ``kernels`` module registers NVRTC-compiled CUDA C++ kernels as
-TensorRT Quick Deployable Plugins. Tensor-only declarative kernels use
-Ahead-of-Time (AOT) plugin launches when available; kernels with
-``ScalarInput`` compile through TensorRT's QDP JIT path because QDP AOT
-extra arguments currently support symbolic integer expressions, not
-arbitrary runtime floats.
-
-A single function — :func:`cuda_kernel_op` — handles both the declarative
-case (drive everything from a :class:`KernelSpec` dataclass) and the
-override case (supply ``meta_fn`` / ``eager_fn`` / ``aot_fn`` / ``schema``
-keyword arguments when the declarative DSL doesn't cover your kernel).
-:func:`ptx_op` is a parallel entry point for kernels that are already
-compiled to PTX bytes.
+The ``kernels`` module registers custom kernels — CUDA C++ compiled with
+NVRTC, or Triton — as TensorRT Quick Deployable Plugins. Tensor-only
+declarative kernels use Ahead-of-Time (AOT) plugin launches when
+available; kernels with ``ScalarInput`` compile through TensorRT's QDP JIT
+path because QDP AOT extra arguments currently support symbolic integer
+expressions, not arbitrary runtime floats.
+
+Three entry points share one registration funnel:
+
+* :func:`cuda_kernel_op` handles both the declarative case (drive
+  everything from a :class:`KernelSpec` dataclass) and the override case
+  (supply ``meta_fn`` / ``eager_fn`` / ``aot_fn`` / ``schema`` keyword
+  arguments when the declarative DSL doesn't cover your kernel).
+* :func:`ptx_op` registers kernels that are already compiled to PTX bytes.
+* :func:`triton_op` registers a ``@triton.jit`` kernel, compiling it to
+  PTX and deriving the AOT launch for you.
 
 Entry points
 ------------
@@ -117,18 +120,97 @@ Pre-compiled PTX entry point
 
 .. autofunction:: ptx_op
 
+Triton entry point
+------------------
+
+.. autofunction:: triton_op
+
+:func:`triton_op` is the Triton analogue of :func:`cuda_kernel_op`. It
+compiles the kernel once with ``triton.compile``, then registers the
+PyTorch custom op, the TRT plugin descriptor, the AOT impl embedding the
+PTX, and the Torch-TensorRT converter — replacing the hand-written
+``@trtp.aot_impl`` boilerplate in the :ref:`aot_plugin` example::
+
+    import tensorrt.plugin as trtp
+    import torch
+    import triton
+    import triton.language as tl
+    import torch_tensorrt.kernels as ttk
+
+    @triton.jit
+    def add_one_kernel(x_ptr, n_elements, y_ptr, BLOCK_SIZE: tl.constexpr):
+        pid = tl.program_id(0)
+        offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
+        mask = offsets < n_elements
+        tl.store(y_ptr + offsets, tl.load(x_ptr + offsets, mask=mask) + 1, mask=mask)
+
+    def add_one_meta(X: torch.Tensor) -> torch.Tensor:
+        return torch.empty_like(X)
+
+    ttk.triton_op(
+        "my::add_one",
+        kernel=add_one_kernel,
+        signature={"x_ptr": "*fp32", "n_elements": "i32", "y_ptr": "*fp32"},
+        constexprs={"BLOCK_SIZE": 256},
+        grid=lambda inputs, outputs: (trtp.cdiv(inputs[0].shape_expr.numel(), 256),),
+        meta_fn=add_one_meta,
+        extra_args_fn=lambda inputs, outputs: [
+            trtp.SymInt32(inputs[0].shape_expr.numel())
+        ],
+    )
+
+``signature`` lists the kernel's non-constexpr parameters in declaration
+order (pointers as ``*<dtype>``, scalars as the bare dtype) and
+``constexprs`` supplies the ``tl.constexpr`` values baked into the PTX.
+``grid`` and ``extra_args_fn`` receive ``trtp.TensorDesc`` objects, so use
+``.shape_expr`` to stay symbolic and keep one engine valid across shapes.
+
+Because the launch is built from symbolic shape expressions, ``triton_op``
+supports dynamic shapes by default.  Pass ``eager_fn`` to also give the op
+a CUDA implementation outside TensorRT, or ``aot_fn`` to replace the
+derived launch entirely.  ``num_warps`` and ``num_stages`` are forwarded to
+``triton.compile``; ``num_warps`` also sets the launch's threads-per-block.
+
+The calling convention is enforced at registration time, because violating
+it does not fail loudly — the kernel would read whatever TensorRT happened
+to place in those argument slots.  :func:`triton_op` raises
+:class:`ValueError` if ``signature`` begins or ends with a scalar,
+interleaves scalars between pointers, disagrees with ``meta_fn``'s arity,
+or declares scalars without an ``extra_args_fn`` to supply them.
+
+.. note::
+
+   A ``triton_op`` registration compiles a single PTX for the given
+   ``signature`` and ``constexprs``.  Inputs or outputs whose dtypes differ
+   from the compiled ones are detected during conversion: the op is left
+   out of the engine and runs in PyTorch, with a warning naming the
+   mismatch.  Register a second op if you need a second dtype —
+   multi-config autotuning and dtype specialization are not yet supported.
+
+   Triton emits PTX at the ISA version of its own bundled ``ptxas``, which
+   can be newer than the installed CUDA driver accepts.  ``triton_op``
+   detects the driver's maximum ISA and recompiles at that version when
+   needed, so the embedded PTX always loads.  It also strips Triton's
+   trailing zero-sized scratch parameters, which TensorRT's AOT launcher
+   does not supply; a kernel needing non-zero scratch cannot use the AOT
+   QDP path and raises at registration time.
+
 Kernel signature convention
 ---------------------------
 
-All entry points assume the ``__global__`` kernel takes its arguments in
-the fixed order::
+All entry points assume the kernel takes its arguments in the fixed
+order::
 
     (input_ptrs..., extras..., output_ptrs...)
 
-Pointers are ``void*`` cast to the appropriate element type.  Extras
-follow the order declared in :attr:`KernelSpec.extras` for the
-declarative path, or the order your ``aot_fn`` builds for the override
-path.
+This matches the order TensorRT passes tensor pointers and AOT extra
+arguments, so no PTX rewriting is needed.  In a CUDA C++ ``__global__``
+kernel, pointers are ``void*`` cast to the appropriate element type; in a
+Triton kernel they are the ``*<dtype>`` parameters declared in
+``signature``.  Extras follow the order declared in
+:attr:`KernelSpec.extras` for the declarative path, the order
+``extra_args_fn`` returns for :func:`triton_op`, or the order your
+``aot_fn`` builds for the override path.
 
 Error behavior
 --------------
diff --git a/docsrc/tutorials/extensibility/plugins/index.rst b/docsrc/tutorials/extensibility/plugins/index.rst
index ef381c0e38..f960883b39 100644
--- a/docsrc/tutorials/extensibility/plugins/index.rst
+++ b/docsrc/tutorials/extensibility/plugins/index.rst
@@ -15,3 +15,4 @@ in serialized engines.
    Example: Custom Kernels with NVRTC in TensorRT AOT Plugins <../../_rendered_examples/dynamo/nvrtc_aot_plugin>
    Example: Auto-derived CUDA Kernel Plugins via cuda_kernel_op <../../_rendered_examples/dynamo/cuda_kernel_op>
    Example: Pre-compiled PTX Kernels via ptx_op <../../_rendered_examples/dynamo/ptx_op>
+   Example: Triton Kernel AOT Plugins via triton_op <../../_rendered_examples/dynamo/triton_op>
diff --git a/docsrc/tutorials/extensibility/plugins/plugins.rst b/docsrc/tutorials/extensibility/plugins/plugins.rst
index 1bfaa22be7..a6c45c17c8 100644
--- a/docsrc/tutorials/extensibility/plugins/plugins.rst
+++ b/docsrc/tutorials/extensibility/plugins/plugins.rst
@@ -23,6 +23,10 @@ depending on your kernel language and performance requirements:
      - Triton
      - Pre-compiled PTX embedded in engine
      - :ref:`aot_plugin`
+   * - QDP declarative (AOT)
+     - Triton
+     - Pre-compiled PTX embedded in engine
+     - :ref:`triton_op`
    * - QDP auto-generate (AOT)
      - CUDA C++ via NVRTC
      - Pre-compiled PTX embedded in engine
@@ -119,6 +123,7 @@ For complete end-to-end examples see:
 
 * :ref:`auto_generate_plugins` — Triton kernel, QDP JIT plugin
 * :ref:`aot_plugin` — Triton kernel, QDP AOT plugin (pre-compiled PTX, no Python overhead at runtime)
+* :ref:`triton_op` — the same Triton AOT plugin registered in one ``triton_op`` call
 * :ref:`nvrtc_aot_plugin` — CUDA C++ kernel compiled with NVRTC, QDP AOT plugin
 * :ref:`custom_kernel_plugins` — manual plugin + converter registration (legacy approach)
 
diff --git a/docsrc/user_guide/compilation/unsupported_ops.rst b/docsrc/user_guide/compilation/unsupported_ops.rst
index 14bac94aee..e9dd194cd9 100644
--- a/docsrc/user_guide/compilation/unsupported_ops.rst
+++ b/docsrc/user_guide/compilation/unsupported_ops.rst
@@ -271,9 +271,10 @@ FAQ
 
 **"I have a Triton kernel. Can I use it in a serialized TRT engine?"**
 
-    Triton kernels can be wrapped as TRT plugins via the AOT plugin path. See
-    :ref:`aot_plugin` for an end-to-end example of compiling a Triton kernel as a
-    TRT plugin for use in a serialized engine.
+    Triton kernels can be wrapped as TRT plugins via the AOT plugin path. The
+    shortest route is ``torch_tensorrt.kernels.triton_op``, which compiles the
+    kernel and registers the op, plugin, and converter in one call — see
+    :ref:`triton_op`. :ref:`aot_plugin` shows the same thing built by hand.
 
 **"Operator X is listed as supported but my model still falls back"**
 
diff --git a/examples/dynamo/triton_op.py b/examples/dynamo/triton_op.py
new file mode 100644
index 0000000000..4496cdff7e
--- /dev/null
+++ b/examples/dynamo/triton_op.py
@@ -0,0 +1,126 @@
+"""
+.. _triton_op:
+
+Register a Triton kernel as an AOT QDP plugin via ``torch_tensorrt.kernels.triton_op``
+======================================================================================
+
+This is the productized form of :ref:`aot_plugin`. That example shows the raw
+mechanism — a ``@triton.jit`` kernel wired to a TensorRT AOT Quick Deployable
+Plugin by hand: you write ``@torch.library.custom_op`` + ``register_fake``,
+``@trtp.register``, and a ~40-line ``@trtp.aot_impl`` that builds the
+``ASTSource`` signature, calls ``triton.compile``, and assembles the launch
+parameters, then ``generate_plugin_converter``.
+
+``triton_op`` collapses all of that into a single call. You provide the Triton
+``signature`` / ``constexprs`` / ``grid`` and a ``meta_fn``; the library
+compiles the kernel to PTX, derives the AOT launch, registers the PyTorch op,
+the plugin descriptor + AOT impl, and the Torch-TensorRT converter.
+
+Calling convention: the kernel's non-constexpr parameters must be declared as
+``(input_ptrs..., extra_scalars..., output_ptrs...)`` so no PTX rewriting is
+needed.
+"""
+
+import argparse
+
+import tensorrt.plugin as trtp
+import torch
+import triton
+import triton.language as tl
+
+import torch_tensorrt
+import torch_tensorrt.kernels as ttk
+
+# %%
+# Step 1: Define the Triton kernel (pure Triton, unchanged from aot_plugin.py)
+# ----------------------------------------------------------------------------
+
+
+@triton.jit
+def add_one_kernel(x_ptr, n_elements, y_ptr, BLOCK_SIZE: tl.constexpr):
+    pid = tl.program_id(0)
+    block_start = pid * BLOCK_SIZE
+    offsets = block_start + tl.arange(0, BLOCK_SIZE)
+    mask = offsets < n_elements
+    x = tl.load(x_ptr + offsets, mask=mask)
+    output = x + 1
+    tl.store(y_ptr + offsets, output, mask=mask)
+
+
+# %%
+# Step 2: Describe the op and register it with a single ``triton_op`` call
+# ------------------------------------------------------------------------
+#
+# * ``meta_fn`` — shape/dtype inference for FakeTensors (the schema is inferred
+#   from its type hints).
+# * ``signature`` / ``constexprs`` — exactly what ``triton.compile`` needs.
+# * ``grid`` / ``extra_args_fn`` — computed from ``trtp.TensorDesc`` inputs,
+#   using symbolic shapes so one engine works across sizes.
+# * ``eager_fn`` — optional; lets ``torch.ops.my.add_one`` also run in eager.
+
+BLOCK_SIZE = 256
+
+
+def add_one_meta(X: torch.Tensor) -> torch.Tensor:
+    return torch.empty_like(X)
+
+
+def add_one_eager(X: torch.Tensor) -> torch.Tensor:
+    Y = torch.empty_like(X)
+    grid = lambda meta: (triton.cdiv(X.numel(), meta["BLOCK_SIZE"]),)
+    add_one_kernel[grid](X, X.numel(), Y, BLOCK_SIZE=BLOCK_SIZE)
+    return Y
+
+
+ttk.triton_op(
+    "my::add_one",
+    kernel=add_one_kernel,
+    signature={"x_ptr": "*fp32", "n_elements": "i32", "y_ptr": "*fp32"},
+    constexprs={"BLOCK_SIZE": BLOCK_SIZE},
+    grid=lambda inputs, outputs: (trtp.cdiv(inputs[0].shape_expr.numel(), BLOCK_SIZE),),
+    meta_fn=add_one_meta,
+    extra_args_fn=lambda inputs, outputs: [trtp.SymInt32(inputs[0].shape_expr.numel())],
+    eager_fn=add_one_eager,
+    supports_dynamic_shapes=True,
+)
+
+
+# %%
+# Step 3: Use it — the op lowers to the AOT QDP plugin inside the engine
+# ----------------------------------------------------------------------
+
+
+class AddOne(torch.nn.Module):
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        return torch.ops.my.add_one(x)
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--min_block_size", type=int, default=1)
+    args = parser.parse_args()
+
+    x = torch.randn(4, 256, device="cuda", dtype=torch.float32)
+    model = AddOne().cuda().eval()
+
+    ref = x + 1
+    eager_out = model(x)
+    assert torch.allclose(eager_out, ref), "eager path mismatch"
+    print("eager path OK")
+
+    trt_model = torch_tensorrt.compile(
+        model,
+        inputs=[x],
+        min_block_size=args.min_block_size,
+    )
+    print("engine compiled with the AOT QDP plugin")
+
+    trt_out = trt_model(x)
+    if torch.allclose(trt_out, ref, atol=1e-5):
+        print("triton_op AOT QDP plugin ran correctly under Torch-TensorRT")
+    else:
+        print(
+            "WARNING: TRT output did not match. Check that the CUDA driver "
+            "supports the kernel's PTX ISA (toolkit newer than driver can "
+            "require capping); registration/compile/build succeeded."
+        )
diff --git a/py/torch_tensorrt/kernels/__init__.py b/py/torch_tensorrt/kernels/__init__.py
index d35e895402..1b56e0424e 100644
--- a/py/torch_tensorrt/kernels/__init__.py
+++ b/py/torch_tensorrt/kernels/__init__.py
@@ -22,6 +22,11 @@
     Useful when the PTX comes from an external compiler (Triton, a cached
     NVRTC output, etc.).
 
+``triton_op`` — register a ``@triton.jit`` kernel. Compiles the kernel to PTX
+    for you and derives the AOT launch, so you only supply the Triton
+    ``signature`` / ``constexprs`` / ``grid`` and a ``meta_fn`` — no
+    hand-written ``@trtp.aot_impl`` compile boilerplate.
+
 Minimal example — declarative ``cuda_kernel_op``::
 
     import torch, torch_tensorrt
@@ -65,7 +70,7 @@
     SameAs,
     ScalarInput,
 )
-from torch_tensorrt.kernels._ops import cuda_kernel_op, ptx_op
+from torch_tensorrt.kernels._ops import cuda_kernel_op, ptx_op, triton_op
 
 __all__ = [
     "Custom",
@@ -81,4 +86,5 @@
     "ScalarInput",
     "cuda_kernel_op",
     "ptx_op",
+    "triton_op",
 ]
diff --git a/py/torch_tensorrt/kernels/_ops.py b/py/torch_tensorrt/kernels/_ops.py
index e81642e0d2..95ffdaf9d8 100644
--- a/py/torch_tensorrt/kernels/_ops.py
+++ b/py/torch_tensorrt/kernels/_ops.py
@@ -1,18 +1,21 @@
 """Public entry points for ``torch_tensorrt.kernels``.
 
-Two functions, two paths into the same registration funnel:
+Three functions, three paths into the same registration funnel:
 
 * :func:`cuda_kernel_op` — declarative entry for CUDA C++ source. Reads
   a :class:`KernelSpec` and derives meta / eager / aot / schema, with
   override keyword arguments for cases outside the DSL.
 * :func:`ptx_op` — escape hatch for pre-compiled PTX bytes (Triton output,
   cached NVRTC artifact). User supplies meta / eager / aot directly.
+* :func:`triton_op` — declarative entry for a ``@triton.jit`` kernel. Compiles
+  the kernel to PTX for you and derives the AOT launch, so callers don't
+  hand-write the ``@trtp.aot_impl`` compile boilerplate.
 """
 
 from __future__ import annotations
 
 import logging
-from typing import Any, Callable, Optional
+from typing import Any, Callable, Dict, Optional
 
 from torch_tensorrt._features import ENABLED_FEATURES
 from torch_tensorrt.dynamo.conversion._ConverterRegistry import ConverterPriority
@@ -26,6 +29,15 @@
 _LOGGER = logging.getLogger(__name__)
 
 
+def _require_qdp_plugin() -> None:
+    """Raise unless the installed TensorRT exposes Quick Deployable Plugins."""
+    if not ENABLED_FEATURES.qdp_plugin:
+        raise RuntimeError(
+            "TensorRT QDP plugins are not available. "
+            "Requires TensorRT >= 10.7.0 (and not 10.14.x)."
+        )
+
+
 def cuda_kernel_op(
     op_name: str,
     spec: KernelSpec,
@@ -61,14 +73,10 @@ def cuda_kernel_op(
     The kernel must follow the calling convention
     ``(input_ptrs..., scalar_inputs..., extras..., output_ptrs...)``.
     """
-    if not ENABLED_FEATURES.qdp_plugin:
-        raise RuntimeError(
-            "TensorRT QDP plugins are not available. "
-            "Requires TensorRT >= 10.7.0 (and not 10.14.x)."
-        )
+    _require_qdp_plugin()
 
     # Late import to avoid circular imports and keep the decorator cheap.
-    from torch_tensorrt.kernels._register import register_cuda_python_plugin
+    from torch_tensorrt.kernels._register import register_qdp_plugin
 
     _validation._validate_spec(
         spec,
@@ -93,7 +101,7 @@ def cuda_kernel_op(
     elif spec.inputs and spec.outputs:
         final_schema = _derive._build_schema(spec)
     else:
-        # Let register_cuda_python_plugin fall back to _infer_schema(meta_fn).
+        # Let register_qdp_plugin fall back to _infer_schema(meta_fn).
         final_schema = None
 
     cuda_spec = CudaPythonSpec(
@@ -120,7 +128,7 @@ def cuda_kernel_op(
             isinstance(input_spec, ScalarInput) for input_spec in (spec.inputs or [])
         )
 
-    register_cuda_python_plugin(
+    register_qdp_plugin(
         op_name=op_name,
         spec=cuda_spec,
         meta_fn=final_meta,
@@ -155,13 +163,9 @@ def ptx_op(
     Use this when the PTX comes from an external compiler (Triton, a cached
     NVRTC output, etc.) and NVRTC compilation should be skipped.
     """
-    if not ENABLED_FEATURES.qdp_plugin:
-        raise RuntimeError(
-            "TensorRT QDP plugins are not available. "
-            "Requires TensorRT >= 10.7.0 (and not 10.14.x)."
-        )
+    _require_qdp_plugin()
 
-    from torch_tensorrt.kernels._register import register_cuda_python_plugin
+    from torch_tensorrt.kernels._register import register_qdp_plugin
 
     spec = CudaPythonSpec(
         kernel_source="",
@@ -169,7 +173,7 @@ def ptx_op(
         aot_fn=aot_fn,
         eager_fn=eager_fn,
     )
-    register_cuda_python_plugin(
+    register_qdp_plugin(
         op_name=op_name,
         spec=spec,
         meta_fn=meta_fn,
@@ -181,3 +185,168 @@ def ptx_op(
         schema=schema,
         precompiled_ptx=ptx,
     )
+
+
+def triton_op(
+    op_name: str,
+    kernel: Any,
+    signature: Dict[str, str],
+    constexprs: Dict[str, Any],
+    grid: Callable[..., Any],
+    meta_fn: Callable[..., Any],
+    *,
+    extra_args_fn: Optional[Callable[..., Any]] = None,
+    aot_fn: Optional[Callable[..., Any]] = None,
+    eager_fn: Optional[Callable[..., Any]] = None,
+    num_warps: Optional[int] = None,
+    num_stages: Optional[int] = None,
+    supports_dynamic_shapes: bool = True,
+    requires_output_allocator: bool = False,
+    priority: ConverterPriority = ConverterPriority.STANDARD,
+    capability_validator: Optional[Callable[..., Any]] = None,
+    schema: Optional[str] = None,
+) -> None:
+    """Register a ``@triton.jit`` kernel as a TensorRT AOT QDP plugin.
+
+    This is the Triton analogue of :func:`cuda_kernel_op`: it compiles the
+    Triton kernel to PTX once (via ``triton.compile``) and wires it through the
+    same registration funnel as ``ptx_op`` — registering the PyTorch custom op,
+    the TRT plugin descriptor, the AOT impl (embedding the PTX), and the
+    Torch-TensorRT converter. It removes the hand-written ``@trtp.aot_impl``
+    compile boilerplate shown in ``examples/dynamo/aot_plugin.py``.
+
+    Calling convention — the Triton kernel's *runtime* parameters (everything
+    except ``tl.constexpr`` args) must be declared in this order::
+
+        (input_ptrs..., extra_scalars..., output_ptrs...)
+
+    and ``signature`` must list those same parameters in the same order. This
+    matches the order TensorRT passes tensor pointers and AOT extra args, so no
+    PTX rewriting is needed.
+
+    Args:
+        op_name: qualified op name ``"ns::name"``. After registration
+            ``torch.ops.ns.name`` exists and is lowered to the QDP plugin
+            during ``torch_tensorrt.compile``.
+        kernel: the ``@triton.jit`` kernel function.
+        signature: Triton signature for the non-constexpr parameters, in
+            declaration order, e.g.
+            ``{"x_ptr": "*fp32", "n_elements": "i32", "y_ptr": "*fp32"}``.
+        constexprs: ``tl.constexpr`` values baked into the PTX,
+            e.g. ``{"BLOCK_SIZE": 256}``.
+        grid: ``callable(inputs, outputs) -> int | tuple`` returning the launch
+            grid, where ``inputs`` / ``outputs`` are ``trtp.TensorDesc`` objects
+            (use ``.shape_expr`` for symbolic dims). Up to three dims are used
+            for ``grid_x`` / ``grid_y`` / ``grid_z``.
+        meta_fn: the fake / meta kernel used for shape+dtype inference. The
+            PyTorch schema is inferred from its type hints unless ``schema`` is
+            passed.
+        extra_args_fn: optional ``callable(inputs, outputs) -> list`` returning
+            the runtime scalar kernel args as ``trtp.SymInt32`` (matching the
+            ``extra_scalars`` in the calling convention). Omit if the kernel
+            has no scalar args.
+        aot_fn: optional full override of the derived AOT launch function
+            (``callable(inputs, outputs, tactic) -> (KernelLaunchParams,
+            extra_args)``). When given, ``grid`` / ``extra_args_fn`` are unused.
+        eager_fn: optional CUDA eager implementation registered on the torch
+            op. Omit if the op is only used through ``torch_tensorrt.compile``.
+        num_warps: warps per block for the compiled kernel, and hence the
+            launch's threads-per-block. Defaults to Triton's own choice.
+        num_stages: software pipelining depth. Defaults to Triton's own choice.
+        capability_validator: optional extra predicate gating conversion. It is
+            combined with the dtype check derived from ``signature`` — both
+            must pass for the op to be lowered to the plugin.
+
+    Raises:
+        ValueError: if ``signature`` does not follow the calling convention, if
+            its pointer counts disagree with ``meta_fn``'s arity, or if it
+            declares scalars without an ``extra_args_fn`` to supply them.
+
+    .. note::
+        This initial implementation compiles a single PTX for the given
+        ``signature`` (fixed input dtypes) and ``constexprs`` (single config).
+        Inputs whose dtypes don't match the compiled ones are declined at
+        conversion time and left to PyTorch. Multi-config autotuning and dtype
+        specialization are follow-up work.
+    """
+    _require_qdp_plugin()
+
+    import tensorrt.plugin as trtp
+
+    from torch_tensorrt.kernels import _triton
+    from torch_tensorrt.kernels._register import register_qdp_plugin, tensor_arity
+    from torch_tensorrt.kernels._triton_spec import TritonSpec
+
+    # Validate before compiling: nothing here needs the kernel built, and every
+    # rule it enforces would otherwise surface as wrong numbers, not an error.
+    layout = _triton.validate_triton_config(
+        op_name,
+        signature,
+        tensor_arity(meta_fn, schema),
+        extra_args_fn,
+        derived_launch=aot_fn is None,
+    )
+
+    ptx, kernel_name, compiled_warps, shared_mem = _triton.compile_triton_to_ptx(
+        kernel, signature, constexprs, num_warps=num_warps, num_stages=num_stages
+    )
+
+    final_validator = _triton.make_dtype_capability_validator(
+        op_name, layout, capability_validator
+    )
+
+    if aot_fn is not None:
+        final_aot = aot_fn
+    else:
+
+        def final_aot(inputs: Any, outputs: Any, tactic: int) -> Any:
+            dims = grid(inputs, outputs)
+            if not isinstance(dims, (tuple, list)):
+                dims = (dims,)
+            if not 1 <= len(dims) <= 3:
+                raise ValueError(
+                    f"triton_op '{op_name}' grid returned {len(dims)} dimension(s); "
+                    "TensorRT launches accept 1 to 3 (grid_x, grid_y, grid_z)."
+                )
+
+            launch_params = trtp.KernelLaunchParams()
+            launch_params.grid_x = dims[0]
+            if len(dims) > 1:
+                launch_params.grid_y = dims[1]
+            if len(dims) > 2:
+                launch_params.grid_z = dims[2]
+            # Triton reports occupancy in warps; TRT wants threads-per-block.
+            launch_params.block_x = compiled_warps * 32
+            launch_params.shared_mem = shared_mem
+
+            if extra_args_fn is None:
+                # The registrar substitutes an empty SymIntExprs for None.
+                return launch_params, None
+
+            values = list(extra_args_fn(inputs, outputs))
+            extra_args = trtp.SymIntExprs(len(values))
+            for idx, value in enumerate(values):
+                extra_args[idx] = value
+            return launch_params, extra_args
+
+    spec = TritonSpec(
+        kernel_name=kernel_name,
+        aot_fn=final_aot,
+        eager_fn=eager_fn,
+        signature=dict(signature),
+        constexprs=dict(constexprs),
+    )
+    register_qdp_plugin(
+        op_name=op_name,
+        spec=spec,
+        meta_fn=meta_fn,
+        supports_dynamic_shapes=supports_dynamic_shapes,
+        requires_output_allocator=requires_output_allocator,
+        priority=priority,
+        capability_validator=final_validator,
+        register_torch_op=True,
+        schema=schema,
+        precompiled_ptx=ptx,
+        use_aot_if_available=True,
+    )
+    _LOGGER.info("triton_op '%s' registered (kernel: %s)", op_name, kernel_name)
diff --git a/py/torch_tensorrt/kernels/_register.py b/py/torch_tensorrt/kernels/_register.py
index bcec7fcdd0..b35da40905 100644
--- a/py/torch_tensorrt/kernels/_register.py
+++ b/py/torch_tensorrt/kernels/_register.py
@@ -2,13 +2,17 @@
 
 import inspect
 import logging
-from typing import Any, Callable, Dict, List, Optional, get_type_hints
+from typing import Any, Callable, Dict, List, Optional, Tuple, Union, get_type_hints
 
 import torch
 
 from torch_tensorrt.dynamo.conversion._ConverterRegistry import ConverterPriority
 from torch_tensorrt.dynamo.conversion.plugins import custom_op
 from torch_tensorrt.kernels._cuda_python_spec import CudaPythonSpec
+from torch_tensorrt.kernels._triton_spec import TritonSpec
+
+# Any kernel spec carrying a compiled entry name + AOT launch function.
+AOTPluginSpec = Union[CudaPythonSpec, TritonSpec]
 
 _LOGGER = logging.getLogger(__name__)
 
@@ -55,33 +59,65 @@ def _patch_trt_shape_expr_reflected_ops() -> None:
 }
 
 
-def _infer_schema(fn: Callable[..., Any]) -> str:
-    """Derive a TorchScript schema like '(Tensor x, int n) -> Tensor' from type hints."""
+def _schema_parts(fn: Callable[..., Any]) -> Tuple[List[Tuple[str, str]], List[str]]:
+    """``([(schema type, arg name)], [return schema types])`` from a fn's hints.
+
+    The single reader of those hints, so a schema built from them and an arity
+    counted from them cannot drift apart.
+    """
     try:
         hints = get_type_hints(fn)
     except Exception:
         hints = {}
 
-    params = list(inspect.signature(fn).parameters.keys())
-    args_str = ", ".join(
-        "{} {}".format(
-            _TORCH_TYPE_TO_SCHEMA.get(hints.get(p, torch.Tensor), "Tensor"), p
-        )
-        for p in params
-    )
+    args = [
+        (_TORCH_TYPE_TO_SCHEMA.get(hints.get(name, torch.Tensor), "Tensor"), name)
+        for name in inspect.signature(fn).parameters
+    ]
 
     ret = hints.get("return", torch.Tensor)
-    origin = getattr(ret, "__origin__", None)
-    if origin is tuple:
-        ret_str = "({})".format(
-            ", ".join(_TORCH_TYPE_TO_SCHEMA.get(t, "Tensor") for t in ret.__args__)
-        )
+    if getattr(ret, "__origin__", None) is tuple:
+        returns = [_TORCH_TYPE_TO_SCHEMA.get(t, "Tensor") for t in ret.__args__]
     else:
-        ret_str = _TORCH_TYPE_TO_SCHEMA.get(ret, "Tensor")
+        returns = [_TORCH_TYPE_TO_SCHEMA.get(ret, "Tensor")]
+    return args, returns
 
+
+def _infer_schema(fn: Callable[..., Any]) -> str:
+    """Derive a TorchScript schema like '(Tensor x, int n) -> Tensor' from type hints."""
+    args, returns = _schema_parts(fn)
+    args_str = ", ".join(f"{schema_type} {name}" for schema_type, name in args)
+    ret_str = returns[0] if len(returns) == 1 else "({})".format(", ".join(returns))
     return f"({args_str}) -> {ret_str}"
 
 
+def tensor_arity(
+    meta_fn: Callable[..., Any], schema: Optional[str] = None
+) -> Optional[Tuple[int, int]]:
+    """``(tensor inputs, outputs)`` of the op that will actually be registered.
+
+    Reads an explicit ``schema`` when one is given, since that is what reaches
+    the dispatcher, and otherwise the same hints :func:`_infer_schema` reads.
+    Returns ``None`` when neither can be interpreted, so callers skip an arity
+    cross-check rather than reject an op over an unreadable annotation.
+    """
+    try:
+        if schema is not None:
+            parsed = torch._C.parse_schema(f"_ttk::_probe{schema}")
+            tensor_type = torch._C.TensorType.get()
+            num_inputs = sum(
+                1 for arg in parsed.arguments if arg.type.isSubtypeOf(tensor_type)
+            )
+            return num_inputs, len(parsed.returns)
+        args, returns = _schema_parts(meta_fn)
+        return sum(1 for schema_type, _ in args if schema_type == "Tensor"), len(
+            returns
+        )
+    except Exception as exc:
+        _LOGGER.debug("Could not determine tensor arity: %s", exc)
+        return None
+
+
 def _torch_op_already_registered(op_name: str) -> bool:
     """Return True if ``op_name`` is already known to the torch dispatcher."""
     ns, name = op_name.split("::", 1)
@@ -142,7 +178,7 @@ def _register_pytorch_op(
     _LOGGER.debug("Registered PyTorch op %s  schema: %s%s", op_name, name, schema_str)
 
 
-def _register_aot_impl(op_name: str, ptx: bytes, spec: CudaPythonSpec) -> None:
+def _register_aot_impl(op_name: str, ptx: bytes, spec: "AOTPluginSpec") -> None:
     """Dynamically build a correctly-typed aot_impl and register it with trtp."""
     from typing import Tuple, Union  # noqa: F401 – used in annotations dict
 
@@ -218,9 +254,9 @@ def _aot_impl({sig}):
     _LOGGER.debug("Registered AOT impl for %s", op_name)
 
 
-def register_cuda_python_plugin(
+def register_qdp_plugin(
     op_name: str,
-    spec: CudaPythonSpec,
+    spec: "AOTPluginSpec",
     meta_fn: Optional[Callable[..., Any]],
     supports_dynamic_shapes: bool = False,
     requires_output_allocator: bool = False,
@@ -231,28 +267,35 @@ def register_cuda_python_plugin(
     precompiled_ptx: Optional[bytes] = None,
     use_aot_if_available: bool = True,
 ) -> None:
-    """Register a NVRTC-compiled CUDA kernel as a TensorRT QDP plugin end-to-end.
+    """Register a kernel (CUDA-C++ or Triton, compiled to PTX) as a TRT QDP plugin.
 
-    Steps performed:
-    1. Compile kernel source to PTX via NVRTC (skipped if ``precompiled_ptx`` is passed).
+    Backend-agnostic: ``spec`` may be a :class:`CudaPythonSpec` (NVRTC source or
+    pre-compiled PTX) or a :class:`TritonSpec` (Triton-compiled PTX). Steps:
+
+    1. Compile kernel source to PTX via NVRTC (skipped if ``precompiled_ptx`` is
+       passed — always the case for Triton).
     2. Optionally register the PyTorch custom op (define + fake impl).
     3. Register the TRT plugin descriptor + JIT impl via generate_plugin().
     4. Register the AOT impl with the compiled PTX.
     5. Register the Torch-TensorRT converter via generate_plugin_converter().
 
-    ``precompiled_ptx`` lets higher-level entry points (e.g. ``cuda_kernel_op``)
-    avoid a redundant second NVRTC pass when they already compiled the source
-    to build an eager kernel handle.
+    ``precompiled_ptx`` lets higher-level entry points (e.g. ``cuda_kernel_op``,
+    ``triton_op``) supply already-compiled PTX and skip the NVRTC pass; only
+    ``CudaPythonSpec`` source-compilation fields are read when it is ``None``.
     """
     if spec.aot_fn is None:
         raise ValueError(
-            f"CudaPythonSpec.aot_fn must be set before registering plugin '{op_name}'. "
-            "Pass aot_fn= to cuda_python() or assign spec.aot_fn directly."
+            f"spec.aot_fn must be set before registering plugin '{op_name}'. "
+            "Pass aot_fn= to the entry point or assign spec.aot_fn directly."
         )
 
     if precompiled_ptx is not None:
         ptx = precompiled_ptx
     else:
+        assert isinstance(spec, CudaPythonSpec), (
+            "source compilation requires a CudaPythonSpec; "
+            "TritonSpec must supply precompiled_ptx"
+        )
         from torch_tensorrt.kernels._nvrtc import compile_to_ptx
 
         ptx, _device, _kernel = compile_to_ptx(
@@ -285,4 +328,4 @@ def register_cuda_python_plugin(
         _aot_register=lambda: _register_aot_impl(op_name, ptx, spec),
     )
 
-    _LOGGER.info("cuda-python QDP plugin '%s' registered successfully", op_name)
+    _LOGGER.info("QDP plugin '%s' registered successfully", op_name)
diff --git a/py/torch_tensorrt/kernels/_triton.py b/py/torch_tensorrt/kernels/_triton.py
new file mode 100644
index 0000000000..d11bba334b
--- /dev/null
+++ b/py/torch_tensorrt/kernels/_triton.py
@@ -0,0 +1,491 @@
+from __future__ import annotations
+
+import functools
+import logging
+import re
+from typing import Any, Callable, Dict, List, NamedTuple, Optional, Tuple
+
+import torch
+
+_LOGGER = logging.getLogger(__name__)
+
+# Triton element-type spellings -> torch dtypes. Spellings absent here (rarer fp8
+# variants, future additions) resolve to None via ``.get``, which disables dtype
+# checking for that parameter rather than rejecting a kernel we don't recognize.
+_TRITON_TO_TORCH_DTYPE: Dict[str, torch.dtype] = {
+    "fp16": torch.float16,
+    "bf16": torch.bfloat16,
+    "fp32": torch.float32,
+    "fp64": torch.float64,
+    "i1": torch.bool,
+    "i8": torch.int8,
+    "i16": torch.int16,
+    "i32": torch.int32,
+    "i64": torch.int64,
+    "u8": torch.uint8,
+    "fp8e4nv": torch.float8_e4m3fn,
+    "fp8e5": torch.float8_e5m2,
+}
+
+
+class SignatureParam(NamedTuple):
+    """One entry of a Triton ``signature``, decoded."""
+
+    name: str
+    is_pointer: bool
+    dtype: Optional[torch.dtype]
+
+
+class SignatureLayout(NamedTuple):
+    """A ``signature`` split along the ``(inputs, extras, outputs)`` convention."""
+
+    inputs: List[SignatureParam]
+    scalars: List[SignatureParam]
+    outputs: List[SignatureParam]
+
+
+def _parse_signature(signature: Dict[str, str]) -> List[SignatureParam]:
+    """Decode each ``signature`` value into pointer-ness and a torch dtype.
+
+    Accepts the ``*fp32`` / ``fp32`` spellings, tolerating Triton's optional
+    ``:<alignment>`` specializer suffix (``*fp32:16``).
+    """
+    params = []
+    for name, type_str in signature.items():
+        text = str(type_str).strip()
+        is_pointer = text.startswith("*")
+        element = text.lstrip("*").split(":", 1)[0].strip()
+        params.append(
+            SignatureParam(name, is_pointer, _TRITON_TO_TORCH_DTYPE.get(element))
+        )
+    return params
+
+
+def analyze_signature(
+    signature: Dict[str, str],
+    arity: Optional[Tuple[int, int]] = None,
+) -> SignatureLayout:
+    """Split a ``signature`` into input pointers, extra scalars, output pointers.
+
+    ``triton_op`` requires the kernel's runtime parameters to be declared as
+    ``(input_ptrs..., extra_scalars..., output_ptrs...)`` because that is the
+    order TensorRT passes tensor pointers and AOT extra arguments. This checks
+    the declaration actually has that shape and, when ``arity`` gives the op's
+    ``(tensor inputs, outputs)`` counts, that the pointer runs are that long.
+
+    Getting this wrong does not fail loudly at runtime — the kernel reads
+    whatever TensorRT happened to place in those slots — so every deviation is
+    rejected here, at registration time.
+
+    Raises:
+        ValueError: if the signature is empty, starts or ends with a scalar,
+            interleaves scalars between pointers, or disagrees with ``arity``.
+    """
+    params = _parse_signature(signature)
+    if not params:
+        raise ValueError(
+            "triton_op signature is empty; it must declare the kernel's "
+            "non-constexpr parameters in declaration order."
+        )
+
+    def convention() -> str:
+        order = ", ".join(
+            f"{p.name}={'ptr' if p.is_pointer else 'scalar'}" for p in params
+        )
+        return (
+            f"Expected (input_ptrs..., extra_scalars..., output_ptrs...); got {order}."
+        )
+
+    if not params[0].is_pointer or not params[-1].is_pointer:
+        raise ValueError(
+            f"triton_op signature must begin and end with pointer parameters. "
+            f"{convention()}"
+        )
+
+    scalar_positions = [i for i, p in enumerate(params) if not p.is_pointer]
+    if scalar_positions and scalar_positions != list(
+        range(scalar_positions[0], scalar_positions[-1] + 1)
+    ):
+        raise ValueError(
+            f"triton_op signature interleaves scalar parameters with pointers. "
+            f"Scalars must form one contiguous run between the input and output "
+            f"pointers. {convention()}"
+        )
+
+    if scalar_positions:
+        # The scalar run pins the boundary, so the split is known either way.
+        num_leading = scalar_positions[0]
+        num_trailing = len(params) - scalar_positions[-1] - 1
+        mismatched = arity is not None and (num_leading, num_trailing) != arity
+    elif arity is not None:
+        # All pointers: only the op's arity can say where inputs end.
+        num_leading, num_trailing = arity
+        mismatched = num_leading + num_trailing != len(params)
+    else:
+        num_leading, num_trailing, mismatched = len(params), 0, False
+
+    if mismatched:
+        assert arity is not None
+        raise ValueError(
+            f"triton_op signature declares {num_leading} leading and "
+            f"{num_trailing} trailing pointer parameter(s) but the op takes "
+            f"{arity[0]} tensor input(s) and returns {arity[1]} output(s). "
+            f"{convention()}"
+        )
+
+    return SignatureLayout(
+        inputs=params[:num_leading],
+        scalars=params[num_leading : len(params) - num_trailing],
+        outputs=params[len(params) - num_trailing :],
+    )
+
+
+def validate_triton_config(
+    op_name: str,
+    signature: Dict[str, str],
+    arity: Optional[Tuple[int, int]],
+    extra_args_fn: Optional[Callable[..., Any]],
+    derived_launch: bool,
+) -> SignatureLayout:
+    """Check a ``triton_op`` registration and return its signature layout.
+
+    Every rule here is answerable from the signature and the op's arity alone,
+    and every one of them, left unchecked, yields wrong numbers rather than an
+    error — so they are enforced before anything is compiled.
+
+    ``derived_launch`` is False when the caller supplied its own ``aot_fn``,
+    which owns the extra arguments and so is exempt from the pairing rules.
+    """
+    layout = analyze_signature(signature, arity)
+    if not derived_launch:
+        return layout
+
+    if layout.scalars and extra_args_fn is None:
+        names = ", ".join(p.name for p in layout.scalars)
+        raise ValueError(
+            f"triton_op '{op_name}' signature declares scalar parameter(s) "
+            f"({names}) but no extra_args_fn was given, so TensorRT would launch "
+            "the kernel with no extra arguments and those scalars would read as "
+            "zero. Pass extra_args_fn=lambda inputs, outputs: [...] returning "
+            "one trtp.SymInt32 per scalar."
+        )
+    if not layout.scalars and extra_args_fn is not None:
+        raise ValueError(
+            f"triton_op '{op_name}' was given an extra_args_fn but its signature "
+            "declares no scalar parameters to receive the values. Add the "
+            "scalars to the signature or drop extra_args_fn."
+        )
+    return layout
+
+
+def make_dtype_capability_validator(
+    op_name: str,
+    layout: SignatureLayout,
+    user_validator: Optional[Callable[..., bool]] = None,
+) -> Callable[..., bool]:
+    """Build a converter capability validator enforcing the compiled dtypes.
+
+    The PTX is compiled once for the dtypes named in ``signature``. Feeding the
+    op tensors of any other dtype reinterprets their bytes and silently returns
+    wrong numbers, so decline the conversion instead: TensorRT then leaves the
+    op to PyTorch rather than embedding a kernel that cannot read its inputs.
+    """
+    expected_inputs = [p.dtype for p in layout.inputs]
+    expected_outputs = [p.dtype for p in layout.outputs]
+
+    def _tensor_meta(value: Any) -> Optional[torch.Tensor]:
+        meta = getattr(value, "meta", None)
+        if not isinstance(meta, dict):
+            return None
+        val = meta.get("val")
+        return val if isinstance(val, torch.Tensor) else None
+
+    def _mismatch(kind: str, index: int, got: torch.Tensor, want: torch.dtype) -> bool:
+        # Warn, not debug: the op silently leaves the engine and runs in
+        # PyTorch, and if no eager_fn was registered the eventual failure is an
+        # opaque "not implemented for the CUDA backend" from the dispatcher.
+        _LOGGER.warning(
+            "Not lowering '%s' to its Triton plugin: %s %d is %s but the kernel "
+            "was compiled for %s. Re-register with a matching signature to run "
+            "it inside TensorRT; it will fall back to PyTorch for now.",
+            op_name,
+            kind,
+            index,
+            got.dtype,
+            want,
+        )
+        return False
+
+    def _validator(node: Any, settings: Any = None) -> bool:
+        if user_validator is not None and not user_validator(node, settings):
+            return False
+
+        actual_inputs = [t for t in map(_tensor_meta, node.args) if t is not None]
+        for index, (got, want) in enumerate(zip(actual_inputs, expected_inputs)):
+            if want is not None and got.dtype != want:
+                return _mismatch("input", index, got, want)
+
+        produced = node.meta.get("val") if isinstance(node.meta, dict) else None
+        actual_outputs = (
+            list(produced) if isinstance(produced, (tuple, list)) else [produced]
+        )
+        for index, (got, want) in enumerate(zip(actual_outputs, expected_outputs)):
+            if want is not None and isinstance(got, torch.Tensor) and got.dtype != want:
+                return _mismatch("output", index, got, want)
+
+        return True
+
+    return _validator
+
+
+def _parse_entry_params(
+    ptx: str, kernel_name: str
+) -> Tuple[Optional[re.Match[str]], List[str]]:
+    """Locate the ``.visible .entry <kernel_name>( ... )`` param list.
+
+    Returns the regex match (groups: prefix, param-block, close-paren) and the
+    list of individual ``.param`` declaration strings, in order.
+    """
+    pattern = re.compile(
+        r"(\.visible\s+\.entry\s+" + re.escape(kernel_name) + r"\s*\()([^)]*)(\))",
+        re.DOTALL,
+    )
+    match = pattern.search(ptx)
+    if match is None:
+        return match, []
+    params = [p.strip() for p in match.group(2).split(",") if p.strip()]
+    return match, params
+
+
+def _strip_trailing_scratch_params(
+    ptx: str,
+    match: Optional[re.Match[str]],
+    params: List[str],
+    keep: int,
+) -> str:
+    """Drop Triton's trailing scratch params so the entry matches TRT's launch.
+
+    Triton (3.x) unconditionally appends ``global_scratch`` and
+    ``profile_scratch`` pointer parameters after the user's kernel arguments.
+    TensorRT's AOT QDP launcher only passes the declared tensor + extra-scalar
+    arguments, so the extra params make the kernel's parameter count exceed what
+    TRT supplies — the launch is misaligned and fails at ``onShapeChange``.
+
+    When those scratch buffers are zero-sized (the common case) the params are
+    unused in the kernel body, so removing them from the ``.entry`` signature is
+    a safe, purely-syntactic fix. ``match`` and ``params`` come from
+    :func:`_parse_entry_params`; ``keep`` is the number of real runtime
+    parameters (i.e. ``len(signature)``).
+    """
+    if match is None or len(params) <= keep:
+        return ptx
+    kept = params[:keep]
+    new_block = "\n\t" + ",\n\t".join(kept) + "\n"
+    return (
+        ptx[: match.start()]
+        + match.group(1)
+        + new_block
+        + match.group(3)
+        + ptx[match.end() :]
+    )
+
+
+@functools.lru_cache(maxsize=1)
+def _driver_max_ptx_version() -> Optional[int]:
+    """Highest PTX ISA the running CUDA driver can load, as a ``ptx_version`` int.
+
+    A driver supports PTX up to the ISA of the CUDA toolkit it ships with, so we
+    read the driver's CUDA version (``cuDriverGetVersion``) and map it with
+    Triton's own ``ptx_get_version`` — the exact mapping Triton uses to decide
+    which ISA to emit. This avoids both a hard-coded version table and trial
+    module loads. Returns ``None`` if it can't be determined (e.g. Triton
+    internals moved), in which case no capping is applied.
+    """
+    try:
+        from cuda.bindings import driver as cuda
+        from triton.backends.nvidia.compiler import ptx_get_version
+
+        cuda.cuInit(0)
+        raw = cuda.cuDriverGetVersion()[1]  # e.g. 13010 -> CUDA 13.1
+        cuda_version = f"{raw // 1000}.{(raw % 1000) // 10}"
+        return int(ptx_get_version(cuda_version))
+    except Exception as exc:  # pragma: no cover - environment dependent
+        _LOGGER.debug("Could not determine driver PTX version: %s", exc)
+        return None
+
+
+@functools.lru_cache(maxsize=1)
+def _triton_default_ptx_version() -> Optional[int]:
+    """The ISA Triton would emit for this GPU, without compiling anything.
+
+    Triton picks its ISA from the version of the ``ptxas`` it bundles, so asking
+    that binary directly predicts the ``.version`` header a compile would
+    produce. Knowing it up front lets the caller compile once at the right ISA
+    instead of compiling, noticing the emitted ISA is too new, and compiling
+    again. Returns ``None`` if Triton's internals moved, which puts the caller
+    back on the compile-then-check path.
+    """
+    try:
+        from triton.backends.nvidia.compiler import get_ptxas, ptx_get_version
+
+        major, minor = torch.cuda.get_device_capability()
+        return int(ptx_get_version(get_ptxas(major * 10 + minor).version))
+    except Exception as exc:  # pragma: no cover - environment dependent
+        _LOGGER.debug("Could not determine Triton's default PTX version: %s", exc)
+        return None
+
+
+def _parse_ptx_version(ptx: str) -> Optional[int]:
+    """``.version 9.3`` -> ``93``, the form Triton's ``ptx_version`` option uses."""
+    match = re.search(r"\.version (\d+)\.(\d+)", ptx)
+    if match is None:
+        return None
+    return int(match.group(1)) * 10 + int(match.group(2))
+
+
+def _triton_import() -> Any:
+    """Import triton, raising an actionable error if it is not installed."""
+    try:
+        import triton
+        import triton.compiler  # noqa: F401  (ensures ASTSource is importable)
+
+        return triton
+    except ImportError:
+        raise ImportError(
+            "triton is required for triton_op plugins. "
+            "Install it with: pip install triton"
+        )
+
+
+def compile_triton_to_ptx(
+    kernel: Any,
+    signature: Dict[str, str],
+    constexprs: Dict[str, Any],
+    num_warps: Optional[int] = None,
+    num_stages: Optional[int] = None,
+) -> Tuple[bytes, str, int, int]:
+    """Compile a ``@triton.jit`` kernel to PTX ahead of time.
+
+    This mirrors what ``examples/dynamo/aot_plugin.py`` does by hand inside its
+    ``@trtp.aot_impl`` body, but performs it once so callers don't have to.
+
+    Args:
+        kernel: the ``@triton.jit`` kernel function.
+        signature: Triton signature mapping *non-constexpr* parameter names to
+            Triton type strings, in kernel-declaration order — e.g.
+            ``{"x_ptr": "*fp32", "n_elements": "i32", "y_ptr": "*fp32"}``.
+            Pointer args use ``*<dtype>``; scalar args use the bare dtype.
+        constexprs: compile-time ``tl.constexpr`` values baked into the PTX,
+            e.g. ``{"BLOCK_SIZE": 256}``. These must NOT appear in ``signature``.
+        num_warps: warps per block, forwarded to ``triton.compile``. Defaults to
+            Triton's own choice. Also sets the launch's threads-per-block.
+        num_stages: software pipelining depth, forwarded to ``triton.compile``.
+            Defaults to Triton's own choice.
+
+    Returns:
+        ``(ptx_bytes, kernel_name, num_warps, shared_mem_bytes)`` — the PTX to
+        embed in the TRT engine, the entry symbol inside it, and the launch
+        metadata needed to build ``trtp.KernelLaunchParams``.
+    """
+    triton = _triton_import()
+
+    base_options: Dict[str, Any] = {}
+    if num_warps is not None:
+        base_options["num_warps"] = num_warps
+    if num_stages is not None:
+        base_options["num_stages"] = num_stages
+
+    def _compile(ptx_version: Optional[int] = None) -> Any:
+        src = triton.compiler.ASTSource(
+            fn=kernel,
+            signature=dict(signature),
+            constexprs=dict(constexprs),
+        )
+        options = dict(base_options)
+        if ptx_version is not None:
+            options["ptx_version"] = ptx_version
+        return triton.compile(src, options=options or None)
+
+    # Triton derives the ISA from its bundled ptxas, which can be newer than the
+    # installed driver; the driver then rejects the PTX at load time with
+    # CUDA_ERROR_UNSUPPORTED_PTX_VERSION, surfacing for AOT plugins as a TRT
+    # ``onShapeChange`` failure at engine runtime. Only lower, never raise — a
+    # ptx_version above what the bundled ptxas can assemble fails the compile.
+    kernel_label = getattr(kernel, "__name__", "<kernel>")
+    driver_max = _driver_max_ptx_version()
+    default_isa = _triton_default_ptx_version()
+    cap = (
+        driver_max
+        if driver_max is not None
+        and default_isa is not None
+        and default_isa > driver_max
+        else None
+    )
+    if cap is not None:
+        _LOGGER.debug(
+            "Triton would emit PTX ISA %d but the driver accepts at most %d; "
+            "compiling '%s' with ptx_version=%d",
+            default_isa,
+            cap,
+            kernel_label,
+            cap,
+        )
+
+    compiled = _compile(ptx_version=cap)
+    ptx_text: str = compiled.asm["ptx"]
+
+    # Fallback for when the ISA couldn't be predicted up front: check what was
+    # actually emitted and pay for a second compile only if it is too new.
+    if cap is None and driver_max is not None:
+        emitted = _parse_ptx_version(ptx_text)
+        if emitted is not None and emitted > driver_max:
+            _LOGGER.debug(
+                "PTX ISA %d exceeds driver max %d; recompiling '%s' with "
+                "ptx_version=%d",
+                emitted,
+                driver_max,
+                kernel_label,
+                driver_max,
+            )
+            compiled = _compile(ptx_version=driver_max)
+            ptx_text = compiled.asm["ptx"]
+
+    kernel_name = compiled.metadata.name
+    # Read back from metadata rather than trusting the request: Triton clamps
+    # num_warps to what the kernel can actually use.
+    compiled_warps = int(compiled.metadata.num_warps)
+    shared_mem = int(compiled.metadata.shared)
+
+    # Zero-sized scratch params are unused and safe to strip; non-zero ones mean
+    # the kernel needs scratch the AOT QDP path cannot provide.
+    global_scratch = int(getattr(compiled.metadata, "global_scratch_size", 0) or 0)
+    profile_scratch = int(getattr(compiled.metadata, "profile_scratch_size", 0) or 0)
+    entry, params = _parse_entry_params(ptx_text, kernel_name)
+    num_runtime_params = len(signature)
+    if len(params) > num_runtime_params:
+        if global_scratch != 0 or profile_scratch != 0:
+            raise RuntimeError(
+                f"Triton kernel '{kernel_name}' requires non-zero scratch memory "
+                f"(global={global_scratch}, profile={profile_scratch}), which the "
+                "TensorRT AOT QDP launch path cannot provide. Rewrite the kernel "
+                "to avoid Triton scratch buffers to use triton_op."
+            )
+        ptx_text = _strip_trailing_scratch_params(
+            ptx_text, entry, params, num_runtime_params
+        )
+        _LOGGER.debug(
+            "Stripped %d trailing scratch param(s) from triton kernel '%s'",
+            len(params) - num_runtime_params,
+            kernel_name,
+        )
+
+    ptx_bytes = ptx_text.encode("utf-8")
+
+    _LOGGER.debug(
+        "Compiled triton kernel '%s' -> PTX (%d bytes, num_warps=%d, shared=%d)",
+        kernel_name,
+        len(ptx_bytes),
+        compiled_warps,
+        shared_mem,
+    )
+    return ptx_bytes, kernel_name, compiled_warps, shared_mem
diff --git a/py/torch_tensorrt/kernels/_triton_spec.py b/py/torch_tensorrt/kernels/_triton_spec.py
new file mode 100644
index 0000000000..c4c039c8cd
--- /dev/null
+++ b/py/torch_tensorrt/kernels/_triton_spec.py
@@ -0,0 +1,26 @@
+from __future__ import annotations
+
+from dataclasses import dataclass, field
+from typing import Any, Callable, Dict, Optional
+
+
+@dataclass(frozen=True)
+class TritonSpec:
+    """Specification for a ``@triton.jit`` kernel registered as an AOT QDP plugin.
+
+    Built by :func:`torch_tensorrt.kernels.triton_op`; not intended to be
+    constructed directly. Carries the compiled entry name, the AOT launch
+    function, an optional CUDA eager implementation, and the Triton
+    ``signature`` / ``constexprs`` for provenance. The compiled PTX itself is
+    passed to the registrar separately as precompiled bytes.
+
+    This is the Triton counterpart to
+    :class:`torch_tensorrt.kernels._cuda_python_spec.CudaPythonSpec`; both are
+    accepted by ``register_qdp_plugin``.
+    """
+
+    kernel_name: str
+    aot_fn: Optional[Callable[..., Any]]
+    eager_fn: Optional[Callable[..., Any]] = None
+    signature: Dict[str, str] = field(default_factory=dict)
+    constexprs: Dict[str, Any] = field(default_factory=dict)
diff --git a/tests/py/kernels/conftest.py b/tests/py/kernels/conftest.py
index 52be6fd4d7..c1ba7c3e4a 100644
--- a/tests/py/kernels/conftest.py
+++ b/tests/py/kernels/conftest.py
@@ -4,6 +4,7 @@
 
 import pytest
 import torch
+
 import torch_tensorrt
 
 skip_no_cuda = pytest.mark.skipif(
@@ -15,20 +16,25 @@
 )
 
 
-def _has_cuda_core() -> bool:
-    """True if the cuda-core ``cuda.core`` API (NVRTC/QDP backend) is importable."""
+def _has_module(*names: str) -> bool:
+    """True if any of ``names`` is importable."""
     import importlib.util
 
-    for mod in ("cuda.core", "cuda.core.experimental"):
+    for name in names:
         try:
-            if importlib.util.find_spec(mod) is not None:
+            if importlib.util.find_spec(name) is not None:
                 return True
         except (ImportError, ModuleNotFoundError, ValueError):
             continue
     return False
 
 
-_HAS_CUDA_CORE = _has_cuda_core()
+skip_no_triton = pytest.mark.skipif(
+    not _has_module("triton"), reason="triton not installed"
+)
+
+# The cuda-core ``cuda.core`` API is the NVRTC/QDP backend.
+_HAS_CUDA_CORE = _has_module("cuda.core", "cuda.core.experimental")
 
 skip_no_cuda_core = pytest.mark.skipif(
     not _HAS_CUDA_CORE,
diff --git a/tests/py/kernels/test_cuda_kernel_op_overrides.py b/tests/py/kernels/test_cuda_kernel_op_overrides.py
index 8e0a800a69..df08446c89 100644
--- a/tests/py/kernels/test_cuda_kernel_op_overrides.py
+++ b/tests/py/kernels/test_cuda_kernel_op_overrides.py
@@ -53,13 +53,13 @@ def meta(x: torch.Tensor, scale: float) -> torch.Tensor:
 
 
 def test_overrides_forward_to_registrar(monkeypatch):
-    """Override kwargs land on register_cuda_python_plugin with the right values."""
+    """Override kwargs land on register_qdp_plugin with the right values."""
     from torch_tensorrt.kernels import _derive, _register
 
     captured = {}
     monkeypatch.setattr(
         _register,
-        "register_cuda_python_plugin",
+        "register_qdp_plugin",
         lambda *a, **k: captured.update(k),
     )
     # Skip the real NVRTC compile — we're testing wiring, not codegen.
@@ -115,7 +115,7 @@ def test_override_missing_required_dsl_field(kwargs, match):
 
 
 def test_precompiled_ptx_skips_nvrtc(monkeypatch):
-    """register_cuda_python_plugin(precompiled_ptx=...) must not call compile_to_ptx."""
+    """register_qdp_plugin(precompiled_ptx=...) must not call compile_to_ptx."""
     from torch_tensorrt.kernels import _nvrtc, _register
     from torch_tensorrt.kernels._cuda_python_spec import CudaPythonSpec
 
@@ -142,7 +142,7 @@ def _fail(*a, **k):
     def _meta(x: torch.Tensor) -> torch.Tensor:
         return torch.empty_like(x)
 
-    _register.register_cuda_python_plugin(
+    _register.register_qdp_plugin(
         op_name="ttk_test::ptx_reused",
         spec=spec,
         meta_fn=_meta,
diff --git a/tests/py/kernels/test_ptx_op.py b/tests/py/kernels/test_ptx_op.py
index 60eff25588..e508ee71a7 100644
--- a/tests/py/kernels/test_ptx_op.py
+++ b/tests/py/kernels/test_ptx_op.py
@@ -23,7 +23,7 @@ def test_ptx_op_forwards_precompiled_ptx(monkeypatch):
     captured = {}
     monkeypatch.setattr(
         _register,
-        "register_cuda_python_plugin",
+        "register_qdp_plugin",
         lambda *a, **k: captured.update(k),
     )
 
@@ -53,7 +53,7 @@ def test_ptx_op_kernel_name_lands_on_spec(monkeypatch):
     captured = {}
     monkeypatch.setattr(
         _register,
-        "register_cuda_python_plugin",
+        "register_qdp_plugin",
         lambda *a, **k: captured.update(k),
     )
 
diff --git a/tests/py/kernels/test_triton_op.py b/tests/py/kernels/test_triton_op.py
new file mode 100644
index 0000000000..55a134801b
--- /dev/null
+++ b/tests/py/kernels/test_triton_op.py
@@ -0,0 +1,555 @@
+"""Tests for triton_op (Triton kernel -> AOT QDP plugin path)."""
+
+import pytest
+import torch
+
+import torch_tensorrt
+import torch_tensorrt.kernels as ttk
+
+from .conftest import (
+    register_once,
+    skip_no_cuda,
+    skip_no_qdp,
+    skip_no_triton,
+)
+
+# A minimal PTX ``.entry`` with two extra trailing (scratch) params, used by the
+# no-GPU unit tests for the PTX post-processing helpers.
+FAKE_PTX = """//
+.version 9.3
+.target sm_90
+.address_size 64
+
+.visible .entry my_kernel(
+\t.param .u64 my_kernel_param_0,
+\t.param .u32 my_kernel_param_1,
+\t.param .u64 my_kernel_param_2,
+\t.param .u64 my_kernel_param_3,
+\t.param .u64 my_kernel_param_4
+)
+{
+\tret;
+}
+"""
+
+
+def _identity_meta(x: torch.Tensor) -> torch.Tensor:
+    return torch.empty_like(x)
+
+
+@pytest.fixture
+def captured_registration(monkeypatch):
+    """Stub compile + registration; returns the kwargs triton_op forwards on."""
+    from torch_tensorrt.kernels import _register, _triton
+
+    monkeypatch.setattr(
+        _triton,
+        "compile_triton_to_ptx",
+        lambda *a, **k: (b"// ptx bytes", "my_kernel_entry", 4, 0),
+    )
+    captured = {}
+    monkeypatch.setattr(
+        _register, "register_qdp_plugin", lambda *a, **k: captured.update(k)
+    )
+    return captured
+
+
+# ---- No-GPU, no-Triton: PTX post-processing helpers ----
+
+
+def test_parse_ptx_version():
+    from torch_tensorrt.kernels._triton import _parse_ptx_version
+
+    assert _parse_ptx_version(FAKE_PTX) == 93
+    assert _parse_ptx_version("// no version here") is None
+
+
+def test_parse_entry_params_counts_all():
+    from torch_tensorrt.kernels._triton import _parse_entry_params
+
+    match, params = _parse_entry_params(FAKE_PTX, "my_kernel")
+    assert match is not None
+    assert len(params) == 5
+
+
+def test_strip_trailing_scratch_params_keeps_only_runtime_args():
+    from torch_tensorrt.kernels._triton import (
+        _parse_entry_params,
+        _strip_trailing_scratch_params,
+    )
+
+    match, params = _parse_entry_params(FAKE_PTX, "my_kernel")
+    stripped = _strip_trailing_scratch_params(FAKE_PTX, match, params, keep=3)
+    _, params = _parse_entry_params(stripped, "my_kernel")
+    assert len(params) == 3
+    assert "my_kernel_param_3" not in stripped
+    assert "my_kernel_param_4" not in stripped
+    # The kept params and body are untouched.
+    assert "my_kernel_param_2" in stripped
+    assert "ret;" in stripped
+
+
+def test_strip_trailing_scratch_params_noop_when_nothing_to_strip():
+    from torch_tensorrt.kernels._triton import (
+        _parse_entry_params,
+        _strip_trailing_scratch_params,
+    )
+
+    match, params = _parse_entry_params(FAKE_PTX, "my_kernel")
+    assert _strip_trailing_scratch_params(FAKE_PTX, match, params, keep=5) == FAKE_PTX
+    # Unknown kernel name -> no match -> left untouched.
+    missing, params = _parse_entry_params(FAKE_PTX, "other")
+    assert _strip_trailing_scratch_params(FAKE_PTX, missing, params, 1) == FAKE_PTX
+
+
+# ---- No-GPU: triton_op plumbing (compile + register mocked out) ----
+
+
+@skip_no_qdp
+def test_triton_op_forwards_to_registrar(captured_registration):
+    """triton_op must compile then forward the PTX + TritonSpec to the registrar."""
+    from torch_tensorrt.kernels._triton_spec import TritonSpec
+
+    captured = captured_registration
+    sig = {"x_ptr": "*fp32", "y_ptr": "*fp32"}
+    ttk.triton_op(
+        "ttk_test::triton_forward",
+        kernel=object(),
+        signature=sig,
+        constexprs={},
+        grid=lambda inputs, outputs: (1,),
+        meta_fn=_identity_meta,
+        supports_dynamic_shapes=True,
+    )
+
+    assert captured["op_name"] == "ttk_test::triton_forward"
+    assert captured["precompiled_ptx"] == b"// ptx bytes"
+    # triton_op builds a TritonSpec, not a CudaPythonSpec.
+    assert isinstance(captured["spec"], TritonSpec)
+    assert captured["spec"].kernel_name == "my_kernel_entry"
+    assert captured["spec"].signature == sig
+    assert captured["use_aot_if_available"] is True
+    assert captured["supports_dynamic_shapes"] is True
+    # A dtype capability validator is always installed, even with none passed.
+    assert callable(captured["capability_validator"])
+
+
+@skip_no_qdp
+def test_triton_op_aot_override_used(captured_registration):
+    """A user-supplied aot_fn overrides the derived one."""
+    captured = captured_registration
+    sentinel = object()
+
+    ttk.triton_op(
+        "ttk_test::triton_override",
+        kernel=object(),
+        signature={"x_ptr": "*fp32", "y_ptr": "*fp32"},
+        constexprs={},
+        grid=lambda i, o: (1,),
+        meta_fn=_identity_meta,
+        aot_fn=lambda *a: sentinel,
+    )
+
+    assert captured["spec"].aot_fn is not None
+    # The override is forwarded verbatim (not wrapped in the derived launcher).
+    assert captured["spec"].aot_fn("i", "o", 0) is sentinel
+
+
+# ---- No-GPU: signature validation (misuse must not reach the GPU) ----
+
+
+def _sig_error(**overrides):
+    """Call triton_op with a deliberately broken config, return the ValueError."""
+    kwargs = dict(
+        kernel=object(),
+        signature={"x_ptr": "*fp32", "y_ptr": "*fp32"},
+        constexprs={},
+        grid=lambda i, o: (1,),
+        meta_fn=_identity_meta,
+    )
+    kwargs.update(overrides)
+    with pytest.raises(ValueError) as excinfo:
+        ttk.triton_op("ttk_test::triton_invalid", **kwargs)
+    return str(excinfo.value)
+
+
+@skip_no_qdp
+def test_scalar_without_extra_args_fn_rejected():
+    """A scalar in the signature with no extra_args_fn would silently read zero."""
+    message = _sig_error(
+        signature={"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32"},
+        extra_args_fn=None,
+    )
+    assert "extra_args_fn" in message
+    assert "n" in message
+
+
+@skip_no_qdp
+def test_extra_args_fn_without_scalars_rejected():
+    message = _sig_error(extra_args_fn=lambda i, o: [1])
+    assert "no scalar parameters" in message
+
+
+@skip_no_qdp
+def test_signature_arity_must_match_meta_fn():
+    """Two tensor inputs + one output needs three pointers, not two."""
+
+    def _meta2(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+        return torch.empty_like(x)
+
+    message = _sig_error(meta_fn=_meta2)
+    assert "2 tensor input(s)" in message
+
+
+@skip_no_qdp
+def test_interleaved_scalar_rejected():
+    message = _sig_error(
+        signature={"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32", "m": "i32"},
+        extra_args_fn=lambda i, o: [1],
+    )
+    assert "begin and end with pointer" in message
+
+
+@skip_no_qdp
+def test_scalar_run_must_be_contiguous():
+    def _meta2(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+        return torch.empty_like(x)
+
+    message = _sig_error(
+        meta_fn=_meta2,
+        signature={
+            "x_ptr": "*fp32",
+            "n": "i32",
+            "y_ptr": "*fp32",
+            "m": "i32",
+            "z_ptr": "*fp32",
+        },
+        extra_args_fn=lambda i, o: [1, 2],
+    )
+    assert "interleaves" in message
+
+
+@skip_no_qdp
+def test_explicit_schema_drives_arity(captured_registration):
+    """An explicit schema, not meta_fn hints, decides what the arity check sees."""
+
+    def _meta(x, n):  # unannotated — hints alone would count both as tensors
+        return torch.empty_like(x)
+
+    ttk.triton_op(
+        "ttk_test::triton_schema_arity",
+        kernel=object(),
+        signature={"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32"},
+        constexprs={},
+        grid=lambda i, o: (1,),
+        meta_fn=_meta,
+        extra_args_fn=lambda i, o: [1],
+        schema="(Tensor x, int n) -> Tensor",
+    )
+    assert captured_registration["schema"] == "(Tensor x, int n) -> Tensor"
+
+
+def test_tensor_arity_prefers_schema_over_hints():
+    from torch_tensorrt.kernels._register import tensor_arity
+
+    def _meta(x, n):
+        return torch.empty_like(x)
+
+    # Unannotated params default to Tensor...
+    assert tensor_arity(_meta) == (2, 1)
+    # ...but an explicit schema is what actually gets registered.
+    assert tensor_arity(_meta, "(Tensor x, int n) -> Tensor") == (1, 1)
+    assert tensor_arity(_meta, "(Tensor x) -> (Tensor, Tensor)") == (1, 2)
+    # An unparseable schema disables the check rather than failing registration.
+    assert tensor_arity(_meta, "not a schema") is None
+
+
+def test_analyze_signature_splits_on_convention():
+    from torch_tensorrt.kernels._triton import analyze_signature
+
+    layout = analyze_signature({"x_ptr": "*fp16", "n": "i32", "y_ptr": "*fp16"}, (1, 1))
+    assert [p.name for p in layout.inputs] == ["x_ptr"]
+    assert [p.name for p in layout.scalars] == ["n"]
+    assert [p.name for p in layout.outputs] == ["y_ptr"]
+    assert layout.inputs[0].dtype == torch.float16
+    assert layout.outputs[0].dtype == torch.float16
+
+
+def test_analyze_signature_tolerates_alignment_suffix_and_unknown_dtype():
+    from torch_tensorrt.kernels._triton import analyze_signature
+
+    layout = analyze_signature({"x_ptr": "*fp32:16", "y_ptr": "*weird"}, (1, 1))
+    assert layout.inputs[0].dtype == torch.float32
+    # Unknown element types disable checking rather than rejecting the kernel.
+    assert layout.outputs[0].dtype is None
+
+
+# ---- No-GPU: dtype capability validator ----
+
+
+class _FakeNode:
+    """Minimal stand-in for the torch.fx.Node a capability validator receives."""
+
+    def __init__(self, arg_dtypes, out_dtype):
+        self.args = [
+            type("_Arg", (), {"meta": {"val": torch.empty(2, dtype=d)}})()
+            for d in arg_dtypes
+        ]
+        self.meta = {"val": torch.empty(2, dtype=out_dtype)}
+
+
+def _validator_for(signature, user_validator=None):
+    from torch_tensorrt.kernels._triton import (
+        analyze_signature,
+        make_dtype_capability_validator,
+    )
+
+    layout = analyze_signature(signature, (1, 1))
+    return make_dtype_capability_validator("ns::op", layout, user_validator)
+
+
+def test_dtype_validator_accepts_matching_dtypes():
+    validate = _validator_for({"x_ptr": "*fp32", "y_ptr": "*fp32"})
+    assert validate(_FakeNode([torch.float32], torch.float32), None) is True
+
+
+def test_dtype_validator_rejects_mismatched_input():
+    """The fp16-into-an-fp32-kernel case, which used to return silent garbage."""
+    validate = _validator_for({"x_ptr": "*fp32", "y_ptr": "*fp32"})
+    assert validate(_FakeNode([torch.float16], torch.float16), None) is False
+
+
+def test_dtype_validator_rejects_mismatched_output():
+    validate = _validator_for({"x_ptr": "*fp32", "y_ptr": "*fp16"})
+    assert validate(_FakeNode([torch.float32], torch.float32), None) is False
+
+
+def test_dtype_validator_composes_with_user_validator():
+    validate = _validator_for(
+        {"x_ptr": "*fp32", "y_ptr": "*fp32"}, user_validator=lambda n, s: False
+    )
+    # dtypes match, but the user's predicate still gets to veto.
+    assert validate(_FakeNode([torch.float32], torch.float32), None) is False
+
+
+def test_dtype_validator_skips_unknown_dtypes():
+    validate = _validator_for({"x_ptr": "*weird", "y_ptr": "*weird"})
+    assert validate(_FakeNode([torch.float16], torch.bfloat16), None) is True
+
+
+# ---- No-GPU: PTX ISA capping ----
+
+# FAKE_PTX declares .version 9.3 and 5 entry params for this 3-param signature.
+CAPPING_SIG = {"a": "*fp32", "b": "i32", "c": "*fp32"}
+
+
+def _stub_triton(monkeypatch, requested, **extra_metadata):
+    """Replace triton.compile with a stub recording each requested ptx_version."""
+    from torch_tensorrt.kernels import _triton
+
+    fake_metadata = type(
+        "_M", (), {"name": "my_kernel", "num_warps": 4, "shared": 0, **extra_metadata}
+    )()
+
+    class _Compiled:
+        asm = {"ptx": FAKE_PTX}
+        metadata = fake_metadata
+
+    class _FakeTriton:
+        class compiler:
+            ASTSource = staticmethod(lambda **kw: object())
+
+        @staticmethod
+        def compile(src, options=None):
+            requested.append((options or {}).get("ptx_version"))
+            return _Compiled()
+
+    monkeypatch.setattr(_triton, "_triton_import", lambda: _FakeTriton)
+
+
+@pytest.mark.parametrize(
+    "driver_max, default_isa, expected",
+    [
+        # Triton's ISA is known up front and too new: one compile, already capped.
+        (91, 93, [91]),
+        # Driver is new enough: one compile at Triton's default.
+        (95, 93, [None]),
+        # Triton's ISA can't be predicted: compile, notice 9.3 > 9.1, recompile.
+        (91, None, [None, 91]),
+        # Unpredictable ISA and unknown driver: no capping at all.
+        (None, None, [None]),
+    ],
+)
+def test_ptx_isa_capped_to_driver(monkeypatch, driver_max, default_isa, expected):
+    from torch_tensorrt.kernels import _triton
+
+    monkeypatch.setattr(_triton, "_driver_max_ptx_version", lambda: driver_max)
+    monkeypatch.setattr(_triton, "_triton_default_ptx_version", lambda: default_isa)
+    requested = []
+    _stub_triton(monkeypatch, requested)
+
+    _triton.compile_triton_to_ptx(object(), CAPPING_SIG, {})
+
+    assert requested == expected
+
+
+def test_nonzero_scratch_raises(monkeypatch):
+    """Triton scratch buffers can't be fed by the AOT launcher — fail loudly."""
+    from torch_tensorrt.kernels import _triton
+
+    monkeypatch.setattr(_triton, "_driver_max_ptx_version", lambda: None)
+    monkeypatch.setattr(_triton, "_triton_default_ptx_version", lambda: None)
+    _stub_triton(monkeypatch, [], global_scratch_size=128, profile_scratch_size=0)
+
+    with pytest.raises(RuntimeError, match="scratch memory"):
+        _triton.compile_triton_to_ptx(object(), CAPPING_SIG, {})
+
+
+@skip_no_qdp
+def test_grid_beyond_three_dims_rejected(captured_registration):
+    """TRT launches take at most grid_x/y/z; extra dims must not be dropped."""
+    captured = captured_registration
+
+    ttk.triton_op(
+        "ttk_test::triton_grid4",
+        kernel=object(),
+        signature={"x_ptr": "*fp32", "y_ptr": "*fp32"},
+        constexprs={},
+        grid=lambda i, o: (1, 2, 3, 4),
+        meta_fn=_identity_meta,
+    )
+
+    with pytest.raises(ValueError, match="4 dimension"):
+        captured["spec"].aot_fn(["in"], ["out"], 0)
+
+
+# ---- GPU integration: real Triton kernel through triton_op ----
+
+try:
+    import triton
+    import triton.language as tl
+
+    @triton.jit
+    def _ttk_add_one_kernel(x_ptr, n, y_ptr, BLOCK: tl.constexpr):
+        pid = tl.program_id(0)
+        off = pid * BLOCK + tl.arange(0, BLOCK)
+        mask = off < n
+        tl.store(y_ptr + off, tl.load(x_ptr + off, mask=mask) + 1, mask=mask)
+
+except ImportError:
+    triton = None
+
+
+def _register_add_one(op_name: str) -> None:
+    import tensorrt.plugin as trtp
+
+    BLOCK = 256
+
+    def _meta(x: torch.Tensor) -> torch.Tensor:
+        return torch.empty_like(x)
+
+    def _eager(x: torch.Tensor) -> torch.Tensor:
+        y = torch.empty_like(x)
+        grid = lambda meta: (triton.cdiv(x.numel(), meta["BLOCK"]),)
+        _ttk_add_one_kernel[grid](x, x.numel(), y, BLOCK=BLOCK)
+        return y
+
+    def _register() -> None:
+        ttk.triton_op(
+            op_name,
+            kernel=_ttk_add_one_kernel,
+            signature={"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32"},
+            constexprs={"BLOCK": BLOCK},
+            grid=lambda inputs, outputs: (
+                trtp.cdiv(inputs[0].shape_expr.numel(), BLOCK),
+            ),
+            meta_fn=_meta,
+            extra_args_fn=lambda inputs, outputs: [
+                trtp.SymInt32(inputs[0].shape_expr.numel())
+            ],
+            eager_fn=_eager,
+            supports_dynamic_shapes=True,
+        )
+
+    register_once(_register)
+
+
+@skip_no_cuda
+@skip_no_qdp
+@skip_no_triton
+class TestTritonOpIntegration:
+    def test_compile_produces_matched_param_count(self):
+        """The compiled+processed PTX entry must have exactly len(signature) params."""
+        from torch_tensorrt.kernels._triton import (
+            _parse_entry_params,
+            compile_triton_to_ptx,
+        )
+
+        sig = {"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32"}
+        ptx, name, num_warps, _shared = compile_triton_to_ptx(
+            _ttk_add_one_kernel, sig, {"BLOCK": 256}
+        )
+        _, params = _parse_entry_params(ptx.decode("utf-8"), name)
+        assert len(params) == len(sig)
+        assert num_warps >= 1
+
+    def test_num_warps_is_honored(self):
+        from torch_tensorrt.kernels._triton import compile_triton_to_ptx
+
+        sig = {"x_ptr": "*fp32", "n": "i32", "y_ptr": "*fp32"}
+        _, _, warps, _ = compile_triton_to_ptx(
+            _ttk_add_one_kernel, sig, {"BLOCK": 256}, num_warps=8, num_stages=2
+        )
+        assert warps == 8
+
+    def test_dtype_mismatch_falls_back_instead_of_returning_garbage(self):
+        """fp16 into an fp32-compiled kernel must not silently produce nonsense."""
+        _register_add_one("ttk_test::triton_add_one_dtype")
+
+        class M(torch.nn.Module):
+            def forward(self, x):
+                return torch.ops.ttk_test.triton_add_one_dtype(x)
+
+        x = torch.randn(4, 256, device="cuda", dtype=torch.float16)
+        trt = torch_tensorrt.compile(
+            M().cuda().eval(),
+            inputs=[x],
+            enabled_precisions={torch.float16},
+            min_block_size=1,
+        )
+        with torch.no_grad():
+            # The plugin is declined, so this runs in PyTorch — and is correct.
+            assert torch.allclose(trt(x), x + 1, atol=1e-2, rtol=1e-2)
+
+    def test_register_and_eager(self):
+        _register_add_one("ttk_test::triton_add_one_eager")
+        x = torch.randn(1024, device="cuda")
+        assert torch.allclose(
+            torch.ops.ttk_test.triton_add_one_eager(x), x + 1, atol=1e-4, rtol=1e-4
+        )
+
+    def test_trt_compile_dynamic_shapes(self):
+        _register_add_one("ttk_test::triton_add_one_dyn")
+
+        class M(torch.nn.Module):
+            def forward(self, x):
+                return torch.ops.ttk_test.triton_add_one_dyn(x)
+
+        inputs = [
+            torch_tensorrt.Input(
+                min_shape=(1, 128),
+                opt_shape=(1, 512),
+                max_shape=(1, 2048),
+                dtype=torch.float32,
+            )
+        ]
+        trt = torch_tensorrt.compile(
+            M().cuda().eval(),
+            inputs=inputs,
+            enabled_precisions={torch.float32},
+            min_block_size=1,
+        )
+        for size in [128, 512, 2048]:
+            x = torch.randn(1, size, device="cuda")
+            with torch.no_grad():
+                assert torch.allclose(trt(x), x + 1, atol=1e-2, rtol=1e-2)