diff --git a/benchmarks/benchmark_rmsnorm.py b/benchmarks/benchmark_rmsnorm.py
new file mode 100644
index 00000000..81325281
--- /dev/null
+++ b/benchmarks/benchmark_rmsnorm.py
@@ -0,0 +1,77 @@
+import argparse
+import time
+
+import torch
+
+from rl_engine.kernels.ops.pytorch.norm.rms_norm import NativeRMSNormOp
+from rl_engine.kernels.ops.triton.rmsnorm_triton import rmsnorm_triton
+
+try:
+    from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE
+    from rl_engine.kernels.ops.cuda.norm.rmsnorm import rmsnorm_cuda
+
+    HAS_CUDA_EXT = _EXT_AVAILABLE and hasattr(_C, "rmsnorm_forward")
+except ImportError:
+    HAS_CUDA_EXT = False
+
+
+def bench(fn, x, w, dy, warmup=20, iters=100):
+    for _ in range(warmup):
+        x.grad = None
+        w.grad = None
+        y = fn(x, w)
+        y.backward(dy)
+    torch.cuda.synchronize()
+
+    start = time.time()
+    for _ in range(iters):
+        x.grad = None
+        w.grad = None
+        y = fn(x, w)
+        y.backward(dy)
+    torch.cuda.synchronize()
+    return (time.time() - start) * 1000.0 / iters
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--T", type=int, default=1024)
+    parser.add_argument("--H", type=int, default=4096)
+    parser.add_argument("--dtype", choices=["fp16", "bf16"], default="bf16")
+    args = parser.parse_args()
+
+    dtype = torch.float16 if args.dtype == "fp16" else torch.bfloat16
+    device = "cuda"
+    T, H = args.T, args.H
+
+    torch.manual_seed(0)
+    x_base = torch.randn((T, H), device=device, dtype=dtype) * 0.2
+    w_base = torch.randn((H,), device=device, dtype=dtype) * 0.2
+    dy = torch.randn((T, H), device=device, dtype=dtype) * 0.2
+
+    def make_inputs():
+        return (
+            x_base.detach().clone().requires_grad_(True),
+            w_base.detach().clone().requires_grad_(True),
+        )
+
+    native = NativeRMSNormOp()
+
+    x, w = make_inputs()
+    t_ref = bench(lambda a, b: native.forward(a, b), x, w, dy)
+    print(f"pytorch ref : {t_ref:.4f} ms")
+
+    x, w = make_inputs()
+    t_tri = bench(lambda a, b: rmsnorm_triton(a, b), x, w, dy)
+    print(f"triton      : {t_tri:.4f} ms | speedup vs ref: {t_ref / t_tri:.2f}x")
+
+    if HAS_CUDA_EXT:
+        x, w = make_inputs()
+        t_cuda = bench(lambda a, b: rmsnorm_cuda(a, b), x, w, dy)
+        print(f"cuda        : {t_cuda:.4f} ms | speedup vs ref: {t_ref / t_cuda:.2f}x")
+    else:
+        print("cuda        : skipped, extension is not built")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/csrc/cuda/rmsnorm.cu b/csrc/cuda/rmsnorm.cu
new file mode 100644
index 00000000..8ab5cde1
--- /dev/null
+++ b/csrc/cuda/rmsnorm.cu
@@ -0,0 +1,327 @@
+#include <torch/extension.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+
+template <typename scalar_t>
+__device__ __forceinline__ float load_as_float(const scalar_t* ptr) {
+    return static_cast<float>(*ptr);
+}
+
+template <>
+__device__ __forceinline__ float load_as_float<at::Half>(const at::Half* ptr) {
+    const __half* p = reinterpret_cast<const __half*>(ptr);
+    return __half2float(*p);
+}
+
+template <>
+__device__ __forceinline__ float load_as_float<at::BFloat16>(const at::BFloat16* ptr) {
+    const __nv_bfloat16* p = reinterpret_cast<const __nv_bfloat16*>(ptr);
+    return __bfloat162float(*p);
+}
+
+
+template <typename scalar_t>
+__device__ __forceinline__ void store_from_float(scalar_t* ptr, float v) {
+    *ptr = static_cast<scalar_t>(v);
+}
+
+template <>
+__device__ __forceinline__ void store_from_float<at::Half>(at::Half* ptr, float v) {
+    __half* p = reinterpret_cast<__half*>(ptr);
+    *p = __float2half(v);
+}
+
+template <>
+__device__ __forceinline__ void store_from_float<at::BFloat16>(at::BFloat16* ptr, float v) {
+    __nv_bfloat16* p = reinterpret_cast<__nv_bfloat16*>(ptr);
+    *p = __float2bfloat16(v);
+}
+
+
+__device__ __forceinline__ float block_reduce_sum(float v) {
+    extern __shared__ float smem[];
+    int tid = threadIdx.x;
+
+    smem[tid] = v;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (tid < stride) {
+            smem[tid] += smem[tid + stride];
+        }
+        __syncthreads();
+    }
+
+    return smem[0];
+}
+
+
+static int choose_threads(int H) {
+    if (H <= 64) return 64;
+    if (H <= 128) return 128;
+    if (H <= 256) return 256;
+    return 512;
+}
+
+constexpr int RMSNORM_DW_ROWS_PER_CHUNK = 256;
+constexpr int RMSNORM_DW_H_TILE = 128;
+
+
+template <typename scalar_t, typename weight_t>
+__global__ void rmsnorm_fwd_kernel(
+    const scalar_t* __restrict__ x,
+    const weight_t* __restrict__ weight,
+    scalar_t* __restrict__ y,
+    float* __restrict__ rstd,
+    int T,
+    int H,
+    float eps
+) {
+    int row = blockIdx.x;
+    int tid = threadIdx.x;
+
+    const scalar_t* x_row = x + row * H;
+    scalar_t* y_row = y + row * H;
+
+    float local_sum = 0.0f;
+
+    // 计算 sum(x^2)，每个 thread 负责若干列。
+    for (int col = tid; col < H; col += blockDim.x) {
+        float xv = load_as_float<scalar_t>(x_row + col);
+        local_sum += xv * xv;
+    }
+
+    // 固定 block reduction。
+    float sum = block_reduce_sum(local_sum);
+
+    float row_rstd = rsqrtf(sum / static_cast<float>(H) + eps);
+
+    if (tid == 0) {
+        rstd[row] = row_rstd;
+    }
+
+    __syncthreads();
+
+    // 写出 y = x * rstd * weight。
+    for (int col = tid; col < H; col += blockDim.x) {
+        float xv = load_as_float<scalar_t>(x_row + col);
+        float wv = load_as_float<weight_t>(weight + col);
+        float out = xv * row_rstd * wv;
+        store_from_float<scalar_t>(y_row + col, out);
+    }
+}
+
+
+template <typename scalar_t, typename weight_t>
+__global__ void rmsnorm_bwd_dx_kernel(
+    const scalar_t* __restrict__ dy,
+    const scalar_t* __restrict__ x,
+    const weight_t* __restrict__ weight,
+    const float* __restrict__ rstd,
+    scalar_t* __restrict__ dx,
+    int T,
+    int H
+) {
+    int row = blockIdx.x;
+    int tid = threadIdx.x;
+
+    const scalar_t* dy_row = dy + row * H;
+    const scalar_t* x_row = x + row * H;
+    scalar_t* dx_row = dx + row * H;
+
+    float local_dot = 0.0f;
+
+    for (int col = tid; col < H; col += blockDim.x) {
+        float dyv = load_as_float<scalar_t>(dy_row + col);
+        float xv = load_as_float<scalar_t>(x_row + col);
+        float wv = load_as_float<weight_t>(weight + col);
+        local_dot += dyv * wv * xv;
+    }
+
+    float dot = block_reduce_sum(local_dot);
+
+    float r = rstd[row];
+    float coeff = dot * r * r * r / static_cast<float>(H);
+
+    for (int col = tid; col < H; col += blockDim.x) {
+        float dyv = load_as_float<scalar_t>(dy_row + col);
+        float xv = load_as_float<scalar_t>(x_row + col);
+        float wv = load_as_float<weight_t>(weight + col);
+
+        float out = r * dyv * wv - xv * coeff;
+        store_from_float<scalar_t>(dx_row + col, out);
+    }
+}
+
+
+template <typename scalar_t>
+__global__ void rmsnorm_partial_dw_kernel(
+    const scalar_t* __restrict__ dy,
+    const scalar_t* __restrict__ x,
+    const float* __restrict__ rstd,
+    const bool* __restrict__ mask,
+    float* __restrict__ partial_dw,
+    int T,
+    int H
+) {
+    int chunk = blockIdx.x;
+    int h = blockIdx.y * RMSNORM_DW_H_TILE + threadIdx.x;
+    if (h >= H) {
+        return;
+    }
+
+    int t0 = chunk * RMSNORM_DW_ROWS_PER_CHUNK;
+    float acc = 0.0f;
+
+#pragma unroll
+    for (int i = 0; i < RMSNORM_DW_ROWS_PER_CHUNK; ++i) {
+        int row = t0 + i;
+        float contrib = 0.0f;
+        if (row < T && mask[row]) {
+            int idx = row * H + h;
+            float dyv = load_as_float<scalar_t>(dy + idx);
+            float xv = load_as_float<scalar_t>(x + idx);
+            contrib = dyv * xv * rstd[row];
+        }
+        acc += contrib;
+    }
+
+    partial_dw[chunk * H + h] = acc;
+}
+
+
+__global__ void rmsnorm_reduce_dw_kernel(
+    const float* __restrict__ partial_dw,
+    float* __restrict__ dw,
+    int chunks,
+    int H
+) {
+    int h = blockIdx.x * RMSNORM_DW_H_TILE + threadIdx.x;
+    if (h >= H) {
+        return;
+    }
+
+    float acc = 0.0f;
+    for (int chunk = 0; chunk < chunks; ++chunk) {
+        acc += partial_dw[chunk * H + h];
+    }
+    dw[h] = acc;
+}
+
+
+void rmsnorm_forward_cuda(
+    torch::Tensor x,
+    torch::Tensor weight,
+    torch::Tensor y,
+    torch::Tensor rstd,
+    double eps
+) {
+    int T = x.size(0);
+    int H = x.size(1);
+    int threads = choose_threads(H);
+    size_t smem = threads * sizeof(float);
+
+    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+    AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, x.scalar_type(), "rmsnorm_forward_cuda", [&] {
+        using x_t = scalar_t;
+        AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, weight.scalar_type(), "rmsnorm_forward_weight_cuda", [&] {
+            using w_t = scalar_t;
+            rmsnorm_fwd_kernel<x_t, w_t><<<T, threads, smem, stream>>>(
+                x.data_ptr<x_t>(),
+                weight.data_ptr<w_t>(),
+                y.data_ptr<x_t>(),
+                rstd.data_ptr<float>(),
+                T,
+                H,
+                static_cast<float>(eps)
+            );
+        });
+    });
+}
+
+
+void rmsnorm_backward_dx_cuda(
+    torch::Tensor dy,
+    torch::Tensor x,
+    torch::Tensor weight,
+    torch::Tensor rstd,
+    torch::Tensor dx
+) {
+    int T = x.size(0);
+    int H = x.size(1);
+    int threads = choose_threads(H);
+    size_t smem = threads * sizeof(float);
+
+    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+    AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, x.scalar_type(), "rmsnorm_backward_dx_cuda", [&] {
+        using x_t = scalar_t;
+        AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, weight.scalar_type(), "rmsnorm_backward_dx_weight_cuda", [&] {
+            using w_t = scalar_t;
+            rmsnorm_bwd_dx_kernel<x_t, w_t><<<T, threads, smem, stream>>>(
+                dy.data_ptr<x_t>(),
+                x.data_ptr<x_t>(),
+                weight.data_ptr<w_t>(),
+                rstd.data_ptr<float>(),
+                dx.data_ptr<x_t>(),
+                T,
+                H
+            );
+        });
+    });
+}
+
+
+void rmsnorm_backward_partial_dw_cuda(
+    torch::Tensor dy,
+    torch::Tensor x,
+    torch::Tensor rstd,
+    torch::Tensor mask,
+    torch::Tensor partial_dw
+) {
+    int T = x.size(0);
+    int H = x.size(1);
+
+    int chunks = (T + RMSNORM_DW_ROWS_PER_CHUNK - 1) / RMSNORM_DW_ROWS_PER_CHUNK;
+    dim3 blocks(chunks, (H + RMSNORM_DW_H_TILE - 1) / RMSNORM_DW_H_TILE);
+    dim3 threads(RMSNORM_DW_H_TILE);
+
+    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+    AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, x.scalar_type(), "rmsnorm_partial_dw_cuda", [&] {
+        rmsnorm_partial_dw_kernel<scalar_t><<<blocks, threads, 0, stream>>>(
+            dy.data_ptr<scalar_t>(),
+            x.data_ptr<scalar_t>(),
+            rstd.data_ptr<float>(),
+            mask.data_ptr<bool>(),
+            partial_dw.data_ptr<float>(),
+            T,
+            H
+        );
+    });
+}
+
+
+void rmsnorm_backward_reduce_dw_cuda(
+    torch::Tensor partial_dw,
+    torch::Tensor dw
+) {
+    int chunks = partial_dw.size(0);
+    int H = partial_dw.size(1);
+
+    dim3 blocks((H + RMSNORM_DW_H_TILE - 1) / RMSNORM_DW_H_TILE);
+    dim3 threads(RMSNORM_DW_H_TILE);
+
+    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+    rmsnorm_reduce_dw_kernel<<<blocks, threads, 0, stream>>>(
+        partial_dw.data_ptr<float>(),
+        dw.data_ptr<float>(),
+        chunks,
+        H
+    );
+}
diff --git a/csrc/ops.cpp b/csrc/ops.cpp
index 61ba4a3b..c6de6aaa 100644
--- a/csrc/ops.cpp
+++ b/csrc/ops.cpp
@@ -75,6 +75,117 @@ torch::Tensor deterministic_logp_forward_fp32(torch::Tensor logits, torch::Tenso
 torch::Tensor deterministic_logp_forward_indexed_out(torch::Tensor logits, torch::Tensor token_ids, torch::Tensor row_indices, torch::Tensor output);
 torch::Tensor deterministic_logp_forward_indexed_fp32(torch::Tensor logits, torch::Tensor token_ids, torch::Tensor row_indices);
 
+// RMSNorm Declarations & Wrappers
+
+constexpr int RMSNORM_DW_ROWS_PER_CHUNK = 256;
+
+void rmsnorm_forward_cuda(
+  torch::Tensor x,
+  torch::Tensor weight,
+  torch::Tensor y,
+  torch::Tensor rstd,
+  double eps);
+
+void rmsnorm_backward_dx_cuda(
+  torch::Tensor dy,
+  torch::Tensor x,
+  torch::Tensor weight,
+  torch::Tensor rstd,
+  torch::Tensor dx);
+
+void rmsnorm_backward_partial_dw_cuda(
+  torch::Tensor dy,
+  torch::Tensor x,
+  torch::Tensor rstd,
+  torch::Tensor mask,
+  torch::Tensor partial_dw);
+
+void rmsnorm_backward_reduce_dw_cuda(
+  torch::Tensor partial_dw,
+  torch::Tensor dw);
+
+static void rmsnorm_check_input(const torch::Tensor& x, const char* name) {
+  TORCH_CHECK(x.is_cuda(), name, " must be a CUDA tensor");
+  TORCH_CHECK(x.is_contiguous(), name, " must be contiguous");
+}
+
+std::vector<torch::Tensor> rmsnorm_forward(
+  torch::Tensor x,
+  torch::Tensor weight,
+  double eps)
+{
+  rmsnorm_check_input(x, "x");
+  rmsnorm_check_input(weight, "weight");
+
+  TORCH_CHECK(x.dim() == 2, "x must be 2D [T, H]");
+  TORCH_CHECK(weight.dim() == 1, "weight must be 1D [H]");
+  TORCH_CHECK(x.size(1) == weight.size(0), "x.size(1) must equal weight.size(0)");
+
+  auto T = x.size(0);
+  auto y = torch::empty_like(x);
+  auto rstd = torch::empty({T}, x.options().dtype(torch::kFloat32));
+
+  rmsnorm_forward_cuda(x, weight, y, rstd, eps);
+
+  return {y, rstd};
+}
+
+torch::Tensor rmsnorm_backward_dx(
+  torch::Tensor dy,
+  torch::Tensor x,
+  torch::Tensor weight,
+  torch::Tensor rstd)
+{
+  rmsnorm_check_input(dy, "dy");
+  rmsnorm_check_input(x, "x");
+  rmsnorm_check_input(weight, "weight");
+  rmsnorm_check_input(rstd, "rstd");
+
+  TORCH_CHECK(dy.sizes() == x.sizes(), "dy and x must have same shape");
+  TORCH_CHECK(x.dim() == 2, "x must be 2D [T, H]");
+  TORCH_CHECK(weight.dim() == 1, "weight must be 1D [H]");
+  TORCH_CHECK(rstd.dim() == 1, "rstd must be 1D [T]");
+  TORCH_CHECK(rstd.size(0) == x.size(0), "rstd.size(0) must equal x.size(0)");
+
+  auto dx = torch::empty_like(x);
+
+  rmsnorm_backward_dx_cuda(dy, x, weight, rstd, dx);
+
+  return dx;
+}
+
+torch::Tensor rmsnorm_backward_dw(
+  torch::Tensor dy,
+  torch::Tensor x,
+  torch::Tensor rstd,
+  torch::Tensor mask)
+{
+  rmsnorm_check_input(dy, "dy");
+  rmsnorm_check_input(x, "x");
+  rmsnorm_check_input(rstd, "rstd");
+  rmsnorm_check_input(mask, "mask");
+
+  TORCH_CHECK(dy.sizes() == x.sizes(), "dy and x must have same shape");
+  TORCH_CHECK(x.dim() == 2, "x must be 2D [T, H]");
+  TORCH_CHECK(rstd.dim() == 1, "rstd must be 1D [T]");
+  TORCH_CHECK(mask.dim() == 1, "mask must be 1D [T]");
+  TORCH_CHECK(mask.scalar_type() == torch::kBool, "mask must be bool");
+  TORCH_CHECK(rstd.size(0) == x.size(0), "rstd.size(0) must equal x.size(0)");
+  TORCH_CHECK(mask.size(0) == x.size(0), "mask.size(0) must equal x.size(0)");
+
+  auto T = x.size(0);
+  auto H = x.size(1);
+  auto chunks = (T + RMSNORM_DW_ROWS_PER_CHUNK - 1) / RMSNORM_DW_ROWS_PER_CHUNK;
+
+  auto partial_dw = torch::empty({chunks, H}, x.options().dtype(torch::kFloat32));
+  auto dw = torch::empty({H}, x.options().dtype(torch::kFloat32));
+
+  rmsnorm_backward_partial_dw_cuda(dy, x, rstd, mask, partial_dw);
+  rmsnorm_backward_reduce_dw_cuda(partial_dw, dw);
+
+  return dw;
+}
+
 // Prefix-Shared Attention Declarations & Wrappers
 
 void prefix_shared_attention_forward(
@@ -169,5 +280,10 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
 
     // registry Prefix-Shared Attention
     m.def("prefix_shared_attention", &prefix_shared_attention, "Prefix-Shared Fused Attention for GRPO");
+
+    // registry RMSNorm
+    m.def("rmsnorm_forward", &rmsnorm_forward, "Batch-invariant RMSNorm forward CUDA");
+    m.def("rmsnorm_backward_dx", &rmsnorm_backward_dx, "Batch-invariant RMSNorm backward dx CUDA");
+    m.def("rmsnorm_backward_dw", &rmsnorm_backward_dw, "Deterministic RMSNorm backward dweight CUDA");
 #endif
 }
diff --git a/rl_engine/kernels/gtest/operator_specs.py b/rl_engine/kernels/gtest/operator_specs.py
index 55a4a203..fff8135a 100644
--- a/rl_engine/kernels/gtest/operator_specs.py
+++ b/rl_engine/kernels/gtest/operator_specs.py
@@ -32,6 +32,19 @@ def _load_object(path: str) -> Any:
 
 
 OP_SPECS = {
+    "rms_norm": OperatorSpec(
+        name="rms_norm",
+        op_class="reduction",
+        gold_path="rl_engine.kernels.ops.pytorch.norm.rms_norm.NativeRMSNormOp",
+        gold_method="forward_fp32",
+        candidate_paths={
+            "pytorch": "rl_engine.kernels.ops.pytorch.norm.rms_norm.NativeRMSNormOp",
+            "triton": "rl_engine.kernels.ops.triton.rmsnorm_triton.RMSNormTritonOp",
+            "cuda": "rl_engine.kernels.ops.cuda.norm.rmsnorm.RMSNormCudaOp",
+            "cuda-sm90": "rl_engine.kernels.ops.cuda.norm.rmsnorm.RMSNormCudaOp",
+        },
+        grad_input_names=("x", "weight"),
+    ),
     "logp": OperatorSpec(
         name="logp",
         op_class="logprob",
diff --git a/rl_engine/kernels/ops/base.py b/rl_engine/kernels/ops/base.py
index 7643e968..ead5c30c 100644
--- a/rl_engine/kernels/ops/base.py
+++ b/rl_engine/kernels/ops/base.py
@@ -3,6 +3,8 @@
 
 from typing import Any
 
+import torch  # noqa: F401  # Load torch shared libraries before importing the binary extension.
+
 from rl_engine.utils.logger import logger
 
 _C: Any = None
diff --git a/rl_engine/kernels/ops/cuda/norm/rmsnorm.py b/rl_engine/kernels/ops/cuda/norm/rmsnorm.py
new file mode 100644
index 00000000..f4adc640
--- /dev/null
+++ b/rl_engine/kernels/ops/cuda/norm/rmsnorm.py
@@ -0,0 +1,89 @@
+import torch
+
+from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE
+
+
+class RMSNormCuda(torch.autograd.Function):
+    """
+    PyTorch autograd wrapper for CUDA RMSNorm.
+    """
+
+    @staticmethod
+    def forward(ctx, x, weight, mask=None, eps=1e-6):
+        """
+        Forward:
+          y = x * rsqrt(mean(x^2) + eps) * weight
+
+        Input:
+          x:      [T, H], fp16/bf16/fp32 CUDA tensor
+          weight: [H],    fp16/bf16/fp32 CUDA tensor
+          mask:   [T],    bool CUDA tensor
+          eps:    float
+
+        Output:
+          y: [T, H]
+        """
+        assert x.is_cuda, "x must be CUDA tensor"
+        assert weight.is_cuda, "weight must be CUDA tensor"
+        assert x.is_contiguous(), "x must be contiguous"
+        assert weight.is_contiguous(), "weight must be contiguous"
+        assert x.dim() == 2, "x must be [T, H]"
+        assert weight.dim() == 1, "weight must be [H]"
+        assert x.shape[1] == weight.shape[0], "hidden size mismatch"
+        assert _EXT_AVAILABLE and hasattr(
+            _C, "rmsnorm_forward"
+        ), "RMSNorm CUDA extension is unavailable. Please rebuild with rmsnorm.cu."
+
+        if mask is None:
+            mask = torch.ones((x.shape[0],), device=x.device, dtype=torch.bool)
+        else:
+            assert mask.is_cuda, "mask must be CUDA tensor"
+            assert mask.is_contiguous(), "mask must be contiguous"
+            assert mask.dtype == torch.bool, "mask must be bool"
+            assert mask.dim() == 1, "mask must be [T]"
+            assert mask.shape[0] == x.shape[0], "mask length mismatch"
+
+        y, rstd = _C.rmsnorm_forward(x, weight, float(eps))
+
+        ctx.save_for_backward(x, weight, rstd, mask)
+        ctx.eps = eps
+
+        return y
+
+    @staticmethod
+    def backward(ctx, grad_out):
+        """
+        Backward:
+          dx = CUDA row-wise deterministic kernel
+          dw = CUDA two-pass deterministic kernel
+        """
+        x, weight, rstd, mask = ctx.saved_tensors
+        dy = grad_out.contiguous()
+
+        dx = _C.rmsnorm_backward_dx(dy, x, weight, rstd)
+
+        dw = _C.rmsnorm_backward_dw(dy, x, rstd, mask)
+
+        return dx, dw, None, None
+
+
+def rmsnorm_cuda(x, weight, eps=1e-6, mask=None):
+    """
+    use:
+        y = rmsnorm_cuda(x, weight)
+        y = rmsnorm_cuda(x, weight, mask=mask)
+    """
+    return RMSNormCuda.apply(x, weight, mask, eps)
+
+
+class RMSNormCudaOp:
+    """CUDA RMSNorm wrapper compatible with the shared operator harness."""
+
+    def __call__(self, x, weight, *, eps=1e-6):
+        return self.forward(x, weight, eps=eps)
+
+    def forward(self, x, weight, *, eps=1e-6):
+        hidden = x.shape[-1]
+        x_2d = x.contiguous().view(-1, hidden)
+        y_2d = rmsnorm_cuda(x_2d, weight.contiguous(), eps=eps)
+        return y_2d.view_as(x)
diff --git a/rl_engine/kernels/ops/triton/rmsnorm_triton.py b/rl_engine/kernels/ops/triton/rmsnorm_triton.py
new file mode 100644
index 00000000..e5979a7c
--- /dev/null
+++ b/rl_engine/kernels/ops/triton/rmsnorm_triton.py
@@ -0,0 +1,111 @@
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _rmsnorm_fwd_kernel(
+    X, W, Y, RSTD, T: tl.constexpr, H: tl.constexpr, EPS: tl.constexpr, BLOCK_H: tl.constexpr
+):
+    row = tl.program_id(0)
+    offs = tl.arange(0, BLOCK_H)
+    mask = offs < H
+
+    x = tl.load(X + row * H + offs, mask=mask, other=0.0).to(tl.float32)
+    w = tl.load(W + offs, mask=mask, other=0.0).to(tl.float32)
+
+    ss = tl.sum(x * x, axis=0)
+    rstd = tl.rsqrt(ss / H + EPS)
+    y = x * rstd * w
+
+    tl.store(Y + row * H + offs, y, mask=mask)
+    tl.store(RSTD + row, rstd)
+
+
+@triton.jit
+def _rmsnorm_bwd_dx_kernel(
+    DY, X, W, RSTD, DX, PARTIAL_DW, T: tl.constexpr, H: tl.constexpr, BLOCK_H: tl.constexpr
+):
+    row = tl.program_id(0)
+    offs = tl.arange(0, BLOCK_H)
+    mask = offs < H
+
+    dy = tl.load(DY + row * H + offs, mask=mask, other=0.0).to(tl.float32)
+    x = tl.load(X + row * H + offs, mask=mask, other=0.0).to(tl.float32)
+    w = tl.load(W + offs, mask=mask, other=0.0).to(tl.float32)
+    rstd = tl.load(RSTD + row).to(tl.float32)
+
+    gw = dy * w
+    dot = tl.sum(gw * x, axis=0)
+    dx = rstd * gw - x * rstd * rstd * rstd * dot / H
+
+    pdw = dy * x * rstd
+
+    tl.store(DX + row * H + offs, dx, mask=mask)
+    tl.store(PARTIAL_DW + row * H + offs, pdw, mask=mask)
+
+
+@triton.jit
+def _rmsnorm_bwd_dw_kernel(PARTIAL_DW, DW, T: tl.constexpr, H: tl.constexpr, BLOCK_T: tl.constexpr):
+    col = tl.program_id(0)
+    offs_t = tl.arange(0, BLOCK_T)
+    mask = offs_t < T
+
+    vals = tl.load(PARTIAL_DW + offs_t * H + col, mask=mask, other=0.0).to(tl.float32)
+    acc = tl.sum(vals, axis=0)
+    tl.store(DW + col, acc)
+
+
+class RMSNormTriton(torch.autograd.Function):
+    @staticmethod
+    def forward(ctx, x, weight, eps: float = 1e-6):
+        assert x.is_cuda and weight.is_cuda
+        assert x.dim() == 2 and weight.dim() == 1
+        T, H = x.shape
+        assert weight.numel() == H
+
+        y = torch.empty_like(x)
+        rstd = torch.empty((T,), device=x.device, dtype=torch.float32)
+
+        block_h = triton.next_power_of_2(H)
+        assert block_h <= 131072, "H too large for this simple Triton kernel"
+
+        _rmsnorm_fwd_kernel[(T,)](x, weight, y, rstd, T, H, eps, BLOCK_H=block_h)
+        ctx.save_for_backward(x, weight, rstd)
+        ctx.H = H
+        return y
+
+    @staticmethod
+    def backward(ctx, grad_out):
+        x, weight, rstd = ctx.saved_tensors
+        T, H = x.shape
+        dx = torch.empty_like(x)
+        partial_dw = torch.empty((T, H), device=x.device, dtype=torch.float32)
+        dw = torch.empty((H,), device=x.device, dtype=torch.float32)
+
+        block_h = triton.next_power_of_2(H)
+        block_t = triton.next_power_of_2(T)
+        assert block_t <= 131072, "T too large for this simple single-program dw reduction"
+
+        _rmsnorm_bwd_dx_kernel[(T,)](
+            grad_out, x, weight, rstd, dx, partial_dw, T, H, BLOCK_H=block_h
+        )
+        _rmsnorm_bwd_dw_kernel[(H,)](partial_dw, dw, T, H, BLOCK_T=block_t)
+        return dx, dw.to(weight.dtype), None
+
+
+def rmsnorm_triton(x, weight, eps: float = 1e-6):
+    return RMSNormTriton.apply(x, weight, eps)
+
+
+class RMSNormTritonOp:
+    """Triton RMSNorm wrapper compatible with the shared operator harness."""
+
+    def __call__(self, x, weight, *, eps: float = 1e-6):
+        return self.forward(x, weight, eps=eps)
+
+    def forward(self, x, weight, *, eps: float = 1e-6):
+        hidden = x.shape[-1]
+        x_2d = x.contiguous().view(-1, hidden)
+        y_2d = rmsnorm_triton(x_2d, weight.contiguous(), eps=eps)
+        return y_2d.view_as(x)
diff --git a/setup.py b/setup.py
index 6f94e040..8d576b7b 100644
--- a/setup.py
+++ b/setup.py
@@ -79,6 +79,7 @@ def get_extensions():
             "csrc/fused_logp_kernel.cu",
             "csrc/deterministic_logp_kernel.cu",
             "csrc/cuda/attention/prefix_shared_attention.cu",
+            "csrc/cuda/rmsnorm.cu",
         ]
 
         cc_major, cc_minor = torch.cuda.get_device_capability()
diff --git a/tests/test_rms_norm.py b/tests/test_rms_norm.py
index d8369abd..5fecc385 100644
--- a/tests/test_rms_norm.py
+++ b/tests/test_rms_norm.py
@@ -5,7 +5,16 @@
 import torch
 import torch.nn.functional as F
 
+from rl_engine.kernels.ops.cuda.norm.rmsnorm import rmsnorm_cuda
 from rl_engine.kernels.ops.pytorch.norm.rms_norm import NativeRMSNormOp
+from rl_engine.kernels.ops.triton.rmsnorm_triton import rmsnorm_triton
+
+try:
+    from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE
+
+    _HAS_CUDA_RMSNORM = _EXT_AVAILABLE and hasattr(_C, "rmsnorm_forward")
+except ImportError:  # pragma: no cover - import can fail when the extension is not built.
+    _HAS_CUDA_RMSNORM = False
 
 # Qwen3-8B normalized dims this op must cover.
 _HIDDEN = 4096  # input / post-attention norm
@@ -26,6 +35,69 @@ def _manual_rms_norm(x, weight, *, eps=_EPS):
     return x_f * torch.rsqrt(var + eps) * weight.float()
 
 
+def _dtype_tolerance(dtype):
+    if dtype is torch.float32:
+        return 2e-5, 2e-5
+    if dtype is torch.float16:
+        return 3e-3, 3e-3
+    if dtype is torch.bfloat16:
+        return 2e-2, 2e-2
+    raise ValueError(f"unsupported dtype: {dtype}")
+
+
+def _run_forward_backward(fn, x, weight, dy):
+    x_req = x.detach().clone().contiguous().requires_grad_(True)
+    w_req = weight.detach().clone().contiguous().requires_grad_(True)
+    y = fn(x_req, w_req)
+    y.backward(dy.detach().clone().contiguous())
+    return y.detach(), x_req.grad.detach(), w_req.grad.detach()
+
+
+def _native_rstd(x, eps=_EPS):
+    return torch.rsqrt(x.float().pow(2).mean(dim=-1) + eps)
+
+
+def _native_dw(x, dy, rstd, mask):
+    contrib = dy.float() * x.float() * rstd.float().unsqueeze(-1)
+    return contrib.masked_fill(~mask[:, None], 0.0).sum(dim=0)
+
+
+def _build_padded_layout(x_real, dy_real, total_rows, real_positions):
+    dtype = x_real.dtype
+    device = x_real.device
+    real_rows, hidden = x_real.shape
+
+    assert len(real_positions) == real_rows
+    assert total_rows >= real_rows
+
+    x_pad = torch.randn((total_rows, hidden), device=device, dtype=torch.float32).to(dtype)
+    dy_pad = torch.randn((total_rows, hidden), device=device, dtype=torch.float32).to(dtype)
+    mask = torch.zeros((total_rows,), device=device, dtype=torch.bool)
+
+    for src_t, dst_t in enumerate(real_positions):
+        x_pad[dst_t] = x_real[src_t]
+        dy_pad[dst_t] = dy_real[src_t]
+        mask[dst_t] = True
+
+    return x_pad, dy_pad, mask
+
+
+def _run_cuda_dw(x, dy, weight, mask, eps=_EPS):
+    x_req = x.detach().clone().contiguous().requires_grad_(True)
+    w_req = weight.detach().clone().contiguous().requires_grad_(True)
+    y = rmsnorm_cuda(x_req, w_req, eps=eps, mask=mask)
+    y.backward(dy.detach().clone().contiguous())
+    return w_req.grad.detach()
+
+
+requires_cuda_rmsnorm = pytest.mark.skipif(
+    not (torch.cuda.is_available() and _HAS_CUDA_RMSNORM),
+    reason="CUDA RMSNorm extension is not available",
+)
+
+requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
+
+
 # 1. Primary correctness check vs PyTorch's own F.rms_norm. This is a *truly*
 # independent implementation -- it may reduce in a different float order than
 # our op, so a shared formula bug (eps placement, wrong reduction dim) cannot
@@ -166,3 +238,151 @@ def test_registry_dispatches_rms_norm():
     op = kernel_registry.get_op("rms_norm")
     assert isinstance(op, NativeRMSNormOp)
     assert hasattr(op, "forward") and hasattr(op, "forward_fp32")
+
+
+@requires_cuda
+@pytest.mark.parametrize("impl", ["triton", "cuda"])
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+@pytest.mark.parametrize("rows, hidden", [(1, 128), (8, 768), (32, 2048), (128, _HIDDEN)])
+def test_cuda_triton_rms_norm_matches_native_forward_and_backward(impl, dtype, rows, hidden):
+    if impl == "cuda" and not _HAS_CUDA_RMSNORM:
+        pytest.skip("CUDA RMSNorm extension is not available")
+
+    torch.manual_seed(0)
+    native = NativeRMSNormOp()
+    x_cpu = torch.randn(rows, hidden, device="cpu", dtype=torch.float32)
+    w_cpu = torch.randn(hidden, device="cpu", dtype=torch.float32)
+    dy_cpu = torch.randn(rows, hidden, device="cpu", dtype=torch.float32)
+
+    x_ref = x_cpu.to(dtype).float().detach().requires_grad_(True)
+    w_ref = w_cpu.to(dtype).float().detach().requires_grad_(True)
+    dy_ref = dy_cpu.to(dtype).float()
+    y_ref = native.forward_fp32(x_ref, w_ref, eps=_EPS)
+    y_ref.backward(dy_ref)
+
+    x_gpu = x_cpu.to(device="cuda", dtype=dtype).detach().requires_grad_(True)
+    w_gpu = w_cpu.to(device="cuda", dtype=dtype).detach().requires_grad_(True)
+    dy_gpu = dy_cpu.to(device="cuda", dtype=dtype)
+    fn = rmsnorm_cuda if impl == "cuda" else rmsnorm_triton
+    y_gpu = fn(x_gpu, w_gpu, eps=_EPS)
+    y_gpu.backward(dy_gpu)
+
+    atol, rtol = _dtype_tolerance(dtype)
+    dw_scale = max(1.0, rows**0.5 / 4.0)
+    torch.testing.assert_close(y_gpu.detach().cpu().float(), y_ref.detach(), atol=atol, rtol=rtol)
+    torch.testing.assert_close(x_gpu.grad.detach().cpu().float(), x_ref.grad, atol=atol, rtol=rtol)
+    torch.testing.assert_close(
+        w_gpu.grad.detach().cpu().float(),
+        w_ref.grad,
+        atol=atol * dw_scale,
+        rtol=rtol * dw_scale,
+    )
+
+
+@requires_cuda
+@pytest.mark.parametrize("impl", ["triton", "cuda"])
+def test_cuda_triton_rms_norm_deterministic_repeat(impl):
+    if impl == "cuda" and not _HAS_CUDA_RMSNORM:
+        pytest.skip("CUDA RMSNorm extension is not available")
+
+    torch.manual_seed(0)
+    fn = rmsnorm_cuda if impl == "cuda" else rmsnorm_triton
+    x = torch.randn(128, _HIDDEN, device="cuda", dtype=torch.bfloat16)
+    weight = torch.randn(_HIDDEN, device="cuda", dtype=torch.bfloat16)
+    dy = torch.randn(128, _HIDDEN, device="cuda", dtype=torch.bfloat16)
+
+    y0, dx0, dw0 = _run_forward_backward(fn, x, weight, dy)
+    torch.cuda.synchronize()
+    for _ in range(10):
+        y, dx, dw = _run_forward_backward(fn, x, weight, dy)
+        torch.cuda.synchronize()
+        assert torch.equal(y0, y)
+        assert torch.equal(dx0, dx)
+        assert torch.equal(dw0, dw)
+
+
+@requires_cuda
+@pytest.mark.parametrize("impl", ["triton", "cuda"])
+@pytest.mark.parametrize("hidden", [63, 64, 65, 127, 128, 129, 255, 256, 257, _HIDDEN])
+def test_cuda_triton_rms_norm_forward_dx_layout_invariance(impl, hidden):
+    if impl == "cuda" and not _HAS_CUDA_RMSNORM:
+        pytest.skip("CUDA RMSNorm extension is not available")
+
+    torch.manual_seed(1)
+    fn = rmsnorm_cuda if impl == "cuda" else rmsnorm_triton
+    dtype = torch.bfloat16
+    target_x = torch.randn(1, hidden, device="cuda", dtype=dtype)
+    target_dy = torch.randn(1, hidden, device="cuda", dtype=dtype)
+    weight = torch.randn(hidden, device="cuda", dtype=dtype)
+
+    y_single, dx_single, _ = _run_forward_backward(fn, target_x, weight, target_dy)
+    for total_rows, row_id in [(16, 0), (16, 7), (64, 63)]:
+        x = torch.randn(total_rows, hidden, device="cuda", dtype=dtype)
+        dy = torch.randn(total_rows, hidden, device="cuda", dtype=dtype)
+        x[row_id : row_id + 1] = target_x
+        dy[row_id : row_id + 1] = target_dy
+        y, dx, _ = _run_forward_backward(fn, x, weight, dy)
+        assert torch.equal(y_single[0], y[row_id])
+        assert torch.equal(dx_single[0], dx[row_id])
+
+    valid_rows = 4
+    positions = [1, 5, 9, 14]
+    valid_x = torch.randn(valid_rows, hidden, device="cuda", dtype=dtype)
+    valid_dy = torch.randn(valid_rows, hidden, device="cuda", dtype=dtype)
+    x_a = torch.randn(16, hidden, device="cuda", dtype=dtype)
+    dy_a = torch.randn(16, hidden, device="cuda", dtype=dtype)
+    x_a[:valid_rows] = valid_x
+    dy_a[:valid_rows] = valid_dy
+    y_a, dx_a, _ = _run_forward_backward(fn, x_a, weight, dy_a)
+
+    x_b = torch.randn(16, hidden, device="cuda", dtype=dtype)
+    dy_b = torch.randn(16, hidden, device="cuda", dtype=dtype)
+    for idx, pos in enumerate(positions):
+        x_b[pos] = valid_x[idx]
+        dy_b[pos] = valid_dy[idx]
+    y_b, dx_b, _ = _run_forward_backward(fn, x_b, weight, dy_b)
+
+    for idx, pos in enumerate(positions):
+        assert torch.equal(y_a[idx], y_b[pos])
+        assert torch.equal(dx_a[idx], dx_b[pos])
+
+
+@requires_cuda_rmsnorm
+def test_cuda_rms_norm_masked_dw_layout_invariance():
+    torch.manual_seed(0)
+    rows, hidden = 128, _HIDDEN
+    dtype = torch.bfloat16
+    x_real = torch.randn((rows, hidden), device="cuda", dtype=torch.float32).to(dtype)
+    dy_real = torch.randn((rows, hidden), device="cuda", dtype=torch.float32).to(dtype)
+    weight = torch.randn((hidden,), device="cuda", dtype=torch.float32).to(dtype)
+
+    x1 = x_real.clone()
+    dy1 = dy_real.clone()
+    mask1 = torch.ones((rows,), device="cuda", dtype=torch.bool)
+    x2, dy2, mask2 = _build_padded_layout(
+        x_real=x_real,
+        dy_real=dy_real,
+        total_rows=2 * rows,
+        real_positions=[2 * i + 1 for i in range(rows)],
+    )
+    x3, dy3, mask3 = _build_padded_layout(
+        x_real=x_real,
+        dy_real=dy_real,
+        total_rows=2 * rows + 1,
+        real_positions=[2 * i for i in range(rows)],
+    )
+
+    dw1 = _run_cuda_dw(x1, dy1, weight, mask1)
+    dw2 = _run_cuda_dw(x2, dy2, weight, mask2)
+    dw3 = _run_cuda_dw(x3, dy3, weight, mask3)
+
+    ref_dw1 = _native_dw(x1, dy1, _native_rstd(x1), mask1)
+    ref_dw2 = _native_dw(x2, dy2, _native_rstd(x2), mask2)
+    ref_dw3 = _native_dw(x3, dy3, _native_rstd(x3), mask3)
+
+    atol, rtol = _dtype_tolerance(dtype)
+    torch.testing.assert_close(dw1.float(), dw2.float(), atol=0.0, rtol=0.0)
+    torch.testing.assert_close(dw2.float(), dw3.float(), atol=0.0, rtol=0.0)
+    torch.testing.assert_close(dw1.float(), ref_dw1.float(), atol=atol, rtol=rtol)
+    torch.testing.assert_close(dw2.float(), ref_dw2.float(), atol=atol, rtol=rtol)
+    torch.testing.assert_close(dw3.float(), ref_dw3.float(), atol=atol, rtol=rtol)