```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

# Custom fused LayerNorm CUDA kernels.
# Strategy: LayerNorm over very large rows (64*256*256 = 4M elements each, 16 rows).
#  - Kernel 1: multi-block-per-row reduction of sum / sum-of-squares with
#    double-precision accumulation + atomicAdd (numerically robust).
#  - Kernel 2: vectorized (float4) normalize+affine pass using precomputed stats.
layernorm_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <ATen/cuda/CUDAContext.h>

#define FULL_MASK 0xffffffffu

__inline__ __device__ double warpReduceSumD(double val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(FULL_MASK, val, offset);
    return val;
}

// Compute per-row sum and sum-of-squares (double precision, atomic accumulation).
__global__ void ln_stats_kernel(const float* __restrict__ x,
                                double* __restrict__ sum,
                                double* __restrict__ sumsq,
                                const int n) {
    const int row = blockIdx.y;
    const float* row_ptr = x + (size_t)row * (size_t)n;

    double s = 0.0;
    double sq = 0.0;

    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int stride = blockDim.x * gridDim.x;

    const int n4 = n >> 2;
    const float4* p4 = reinterpret_cast<const float4*>(row_ptr);
    for (int i = tid; i < n4; i += stride) {
        float4 v = p4[i];
        s  += (double)v.x + (double)v.y + (double)v.z + (double)v.w;
        sq += (double)v.x * v.x + (double)v.y * v.y +
              (double)v.z * v.z + (double)v.w * v.w;
    }
    for (int i = (n4 << 2) + tid; i < n; i += stride) {
        float v = row_ptr[i];
        s += (double)v;
        sq += (double)v * v;
    }

    // Block reduction.
    s = warpReduceSumD(s);
    sq = warpReduceSumD(sq);

    __shared__ double s_sum[32];
    __shared__ double s_sq[32];
    const int lane = threadIdx.x & 31;
    const int wid = threadIdx.x >> 5;
    if (lane == 0) {
        s_sum[wid] = s;
        s_sq[wid] = sq;
    }
    __syncthreads();
    const int nwarps = (blockDim.x + 31) >> 5;
    if (wid == 0) {
        double vs = (lane < nwarps) ? s_sum[lane] : 0.0;
        double vq = (lane < nwarps) ? s_sq[lane] : 0.0;
        vs = warpReduceSumD(vs);
        vq = warpReduceSumD(vq);
        if (lane == 0) {
            atomicAdd(&sum[row], vs);
            atomicAdd(&sumsq[row], vq);
        }
    }
}

// Normalize + affine transform, vectorized with float4.
__global__ void ln_apply_kernel(const float* __restrict__ x,
                                const float* __restrict__ w,
                                const float* __restrict__ b,
                                const double* __restrict__ sum,
                                const double* __restrict__ sumsq,
                                float* __restrict__ y,
                                const int n,
                                const float eps) {
    const int row = blockIdx.y;
    const double mean = sum[row] / (double)n;
    double var = sumsq[row] / (double)n - mean * mean;
    var = var > 0.0 ? var : 0.0;
    const float rstd = rsqrtf((float)var + eps);
    const float mean_f = (float)mean;

    const float* row_x = x + (size_t)row * (size_t)n;
    float* row_y = y + (size_t)row * (size_t)n;

    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int stride = blockDim.x * gridDim.x;

    const int n4 = n >> 2;
    const float4* x4 = reinterpret_cast<const float4*>(row_x);
    const float4* w4 = reinterpret_cast<const float4*>(w);
    const float4* b4 = reinterpret_cast<const float4*>(b);
    float4* y4 = reinterpret_cast<float4*>(row_y);
    for (int i = tid; i < n4; i += stride) {
        float4 xv = x4[i];
        float4 wv = w4[i];
        float4 bv = b4[i];
        float4 yv;
        // y = x * (rstd*w) + (b - mean*rstd*w)  -> one FMA per element after 2 setup ops
        float r0 = rstd * wv.x;
        float r1 = rstd * wv.y;
        float r2 = rstd * wv.z;
        float r3 = rstd * wv.w;
        yv.x = fmaf(xv.x, r0, fmaf(-mean_f, r0, bv.x));
        yv.y = fmaf(xv.y, r1, fmaf(-mean_f, r1, bv.y));
        yv.z = fmaf(xv.z, r2, fmaf(-mean_f, r2, bv.z));
        yv.w = fmaf(xv.w, r3, fmaf(-mean_f, r3, bv.w));
        y4[i] = yv;
    }
    for (int i = (n4 << 2) + tid; i < n; i += stride) {
        row_y[i] = (row_x[i] - mean_f) * rstd * w[i] + b[i];
    }
}

torch::Tensor layernorm_forward_cuda(torch::Tensor x,
                                     torch::Tensor weight,
                                     torch::Tensor bias,
                                     double eps,
                                     int64_t normalized_numel) {
    TORCH_CHECK(x.is_cuda(), "x must be CUDA");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");

    auto x_c = x.contiguous();
    auto w_c = weight.contiguous();
    auto b_c = bias.contiguous();

    const int n = (int)normalized_numel;
    const int rows = (int)(x_c.numel() / normalized_numel);

    auto y = torch::empty_like(x_c);
    auto opts = torch::TensorOptions().dtype(torch::kFloat64).device(x.device());
    auto stats = torch::zeros({2, rows}, opts);
    double* sum_ptr = stats.data_ptr<double>();
    double* sumsq_ptr = sum_ptr + rows;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const int block = 256;
    int blocks_per_row = 64;
    const int max_blocks_x = (n / 4 + block - 1) / block;
    if (blocks_per_row > max_blocks_x) blocks_per_row = max_blocks_x;
    if (blocks_per_row < 1) blocks_per_row = 1;
    dim3 grid(blocks_per_row, rows);

    ln_stats_kernel<<<grid, block, 0, stream>>>(
        x_c.data_ptr<float>(), sum_ptr, sumsq_ptr, n);
    ln_apply_kernel<<<grid, block, 0, stream>>>(
        x_c.data_ptr<float>(), w_c.data_ptr<float>(), b_c.data_ptr<float>(),
        sum_ptr, sumsq_ptr, y.data_ptr<float>(), n, (float)eps);

    return y;
}
"""

layernorm_cpp_source = (
    "torch::Tensor layernorm_forward_cuda(torch::Tensor x, torch::Tensor weight, "
    "torch::Tensor bias, double eps, int64_t normalized_numel);"
)

_layernorm_ext = load_inline(
    name="custom_fused_layernorm",
    cpp_sources=layernorm_cpp_source,
    cuda_sources=layernorm_cuda_source,
    functions=["layernorm_forward_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    LayerNorm model accelerated with a custom fused CUDA kernel.
    Keeps the same parameter structure (self.ln) as the reference model.
    """
    def __init__(self, normalized_shape: tuple):
        super(ModelNew, self).__init__()
        self.ln = nn.LayerNorm(normalized_shape=normalized_shape)
        self.normalized_shape = tuple(self.ln.normalized_shape)
        self.normalized_numel = 1
        for d in self.normalized_shape:
            self.normalized_numel *= int(d)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if (x.is_cuda and x.dtype == torch.float32
                and self.ln.elementwise_affine):
            return _layernorm_ext.layernorm_forward_cuda(
                x, self.ln.weight, self.ln.bias,
                float(self.ln.eps), self.normalized_numel)
        # Robust fallback for CPU tensors / other dtypes.
        return F.layer_norm(x, self.ln.normalized_shape,
                            self.ln.weight, self.ln.bias, self.ln.eps)
```