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

_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* shared) {
    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)
        shared[warp] = value;
    __syncthreads();

    if (warp == 0) {
        value = lane < 8 ? shared[lane] : 0.0f;

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

        if (lane == 0)
            shared[8] = value;
    }
    __syncthreads();
    return shared[8];
}

__global__ void norm_residual_mul_kernel(
    const float* __restrict__ input,
    const float* __restrict__ residual,
    float* __restrict__ output,
    float eps
) {
    constexpr int WIDTH = 8192;
    constexpr int THREADS = 256;
    constexpr int VECTORS = WIDTH / (THREADS * 4);

    __shared__ float reduction[9];

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

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

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

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

    const float mean = block_sum(sum, reduction) / WIDTH;

    float squared_sum = 0.0f;
    #pragma unroll
    for (int i = 0; i < VECTORS; ++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;
        squared_sum += v.y * v.y;
        squared_sum += v.z * v.z;
        squared_sum += v.w * v.w;
    }

    const float variance = block_sum(squared_sum, reduction) / WIDTH;
    const float inverse_std = rsqrtf(variance + eps);

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

        float4 result;
        result.x = (__fmul_rn(v.x, inverse_std) + y.x) * y.x;
        result.y = (__fmul_rn(v.y, inverse_std) + y.y) * y.y;
        result.z = (__fmul_rn(v.z, inverse_std) + y.z) * y.z;
        result.w = (__fmul_rn(v.w, inverse_std) + y.w) * y.w;
        output4[index] = result;
    }
}

torch::Tensor norm_residual_mul(
    torch::Tensor input,
    torch::Tensor residual,
    double eps
) {
    TORCH_CHECK(input.is_cuda() && residual.is_cuda(),
                "Inputs must be CUDA tensors");
    TORCH_CHECK(input.device() == residual.device(),
                "Inputs must be on the same device");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                residual.scalar_type() == at::kFloat,
                "Inputs must be float32");
    TORCH_CHECK(input.dim() == 2 && input.size(1) == 8192,
                "Input must have shape (batch, 8192)");
    TORCH_CHECK(input.sizes() == residual.sizes(),
                "Input shapes must match");
    TORCH_CHECK(input.is_contiguous() && residual.is_contiguous(),
                "Inputs must be contiguous");

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

    if (input.size(0) > 0) {
        norm_residual_mul_kernel<<<
            input.size(0), 256, 0,
            at::cuda::getCurrentCUDAStream()
        >>>(
            input.data_ptr<float>(),
            residual.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_mul_8192",
    cpp_sources=(
        "torch::Tensor norm_residual_mul("
        "torch::Tensor input, torch::Tensor residual, double eps);"
    ),
    cuda_sources=_cuda_source,
    functions=["norm_residual_mul"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    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):
        x = self.bmm(x)

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

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