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


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor fused_norm_residual_cuda(
    torch::Tensor x, torch::Tensor y, double eps);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float block_sum(float value, float* scratch) {
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;

    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        value += __shfl_down_sync(0xffffffff, value, offset);

    if (lane == 0)
        scratch[warp] = value;
    __syncthreads();

    if (warp == 0) {
        value = lane < 8 ? scratch[lane] : 0.0f;
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1)
            value += __shfl_down_sync(0xffffffff, value, offset);
        if (lane == 0)
            scratch[0] = value;
    }
    __syncthreads();

    const float result = scratch[0];
    __syncthreads();
    return result;
}

__global__ void fused_norm_residual_8192(
    const float* __restrict__ x,
    const float* __restrict__ y,
    float* __restrict__ output,
    float eps
) {
    __shared__ float scratch[8];

    const int tid = threadIdx.x;
    const int64_t row_offset = static_cast<int64_t>(blockIdx.x) * 8192;

    const float4* input4 =
        reinterpret_cast<const float4*>(x + row_offset);
    const float4* residual4 =
        reinterpret_cast<const float4*>(y + row_offset);
    float4* output4 =
        reinterpret_cast<float4*>(output + row_offset);

    float4 values[8];
    float sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        const float4 v = input4[tid + i * 256];
        values[i] = v;
        sum += (v.x + v.y) + (v.z + v.w);
    }

    const float mean = block_sum(sum, scratch) * (1.0f / 8192.0f);
    float squared_sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        float4 v = values[i];
        v.x -= mean;
        v.y -= mean;
        v.z -= mean;
        v.w -= mean;
        values[i] = v;
        squared_sum +=
            (v.x * v.x + v.y * v.y) +
            (v.z * v.z + v.w * v.w);
    }

    const float variance =
        block_sum(squared_sum, scratch) * (1.0f / 8192.0f);
    const float inverse_std = rsqrtf(variance + eps);

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        const int index = tid + i * 256;
        const float4 v = values[i];
        const float4 r = residual4[index];

        output4[index] = make_float4(
            (v.x * inverse_std + r.x) * r.x,
            (v.y * inverse_std + r.y) * r.y,
            (v.z * inverse_std + r.z) * r.z,
            (v.w * inverse_std + r.w) * r.w
        );
    }
}

torch::Tensor fused_norm_residual_cuda(
    torch::Tensor x, torch::Tensor y, double eps
) {
    TORCH_CHECK(x.is_cuda() && y.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.device() == y.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32 &&
                y.scalar_type() == torch::kFloat32,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 2 && x.size(1) == 8192,
                "Expected shape [batch, 8192]");
    TORCH_CHECK(x.sizes() == y.sizes(), "Shape mismatch");
    TORCH_CHECK(x.is_contiguous() && y.is_contiguous(),
                "Expected contiguous tensors");

    const c10::cuda::CUDAGuard guard(x.device());
    auto output = torch::empty_like(x);

    if (x.size(0) != 0) {
        fused_norm_residual_8192<<<
            static_cast<unsigned int>(x.size(0)),
            256,
            0,
            at::cuda::getCurrentCUDAStream()
        >>>(
            x.data_ptr<float>(),
            y.data_ptr<float>(),
            output.data_ptr<float>(),
            static_cast<float>(eps)
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_fused_ops = load_inline(
    name="fused_instance_norm_residual_8192_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_norm_residual_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self, in_features, out_features, eps=1e-5, momentum=0.1):
        super().__init__()
        self.bmm = nn.Linear(in_features, out_features)
        self.instance_norm = nn.InstanceNorm2d(
            out_features, eps=eps, momentum=momentum
        )

    def forward(self, x, y):
        z = F.linear(x, self.bmm.weight, self.bmm.bias)

        if (
            z.is_cuda
            and y.is_cuda
            and z.device == y.device
            and z.dtype == torch.float32
            and y.dtype == torch.float32
            and z.ndim == 2
            and z.shape[1] == 8192
            and z.shape == y.shape
            and not (torch.is_grad_enabled() and (z.requires_grad or y.requires_grad))
        ):
            return _fused_ops.fused_norm_residual_cuda(
                z.contiguous(), y.contiguous(), self.instance_norm.eps
            )

        z = self.instance_norm(
            z.unsqueeze(1).unsqueeze(1)
        ).squeeze(1).squeeze(1)
        return (z + y) * y
```