```python
import torch
import torch.nn as nn
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>
#include <cstdint>

constexpr int THREADS = 256;

__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();

    value = lane < THREADS / 32 ? scratch[lane] : 0.0f;

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

    // Ensure every warp has finished reading before scratch is reused.
    __syncthreads();
    return __shfl_sync(0xffffffff, value, 0);
}

__global__ void fused_norm_8192(
    const float* __restrict__ x,
    const float* __restrict__ y,
    float* __restrict__ output,
    float eps
) {
    constexpr int WIDTH = 8192;
    constexpr int ITEMS = WIDTH / (THREADS * 4);

    __shared__ float scratch[THREADS / 32];

    const int64_t row_offset = static_cast<int64_t>(blockIdx.x) * WIDTH;
    const float4* row_x = reinterpret_cast<const float4*>(x + row_offset);
    const float4* row_y = reinterpret_cast<const float4*>(y + row_offset);
    float4* row_output = reinterpret_cast<float4*>(output + row_offset);

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

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        float4 value = row_x[threadIdx.x + i * THREADS];
        values[i] = value;
        sum += (value.x + value.y) + (value.z + value.w);
    }

    const float mean = block_sum(sum, scratch) / WIDTH;
    float squared_sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        float4 value = values[i];
        value.x -= mean;
        value.y -= mean;
        value.z -= mean;
        value.w -= mean;
        values[i] = value;

        squared_sum +=
            (value.x * value.x + value.y * value.y) +
            (value.z * value.z + value.w * value.w);
    }

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

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int index = threadIdx.x + i * THREADS;
        const float4 residual = row_y[index];
        const float4 value = values[i];

        float4 result;
        result.x = (value.x * inverse_std + residual.x) * residual.x;
        result.y = (value.y * inverse_std + residual.y) * residual.y;
        result.z = (value.z * inverse_std + residual.z) * residual.z;
        result.w = (value.w * inverse_std + residual.w) * residual.w;
        row_output[index] = result;
    }
}

__global__ void fused_norm_generic(
    const float* __restrict__ x,
    const float* __restrict__ y,
    float* __restrict__ output,
    int64_t width,
    float eps
) {
    __shared__ float scratch[THREADS / 32];

    const int64_t row_offset = static_cast<int64_t>(blockIdx.x) * width;
    x += row_offset;
    y += row_offset;
    output += row_offset;

    float sum = 0.0f;
    for (int64_t col = threadIdx.x; col < width; col += THREADS)
        sum += x[col];

    const float mean = block_sum(sum, scratch) / static_cast<float>(width);

    float squared_sum = 0.0f;
    for (int64_t col = threadIdx.x; col < width; col += THREADS) {
        const float centered = x[col] - mean;
        squared_sum += centered * centered;
    }

    const float variance =
        block_sum(squared_sum, scratch) / static_cast<float>(width);
    const float inverse_std = rsqrtf(variance + eps);

    for (int64_t col = threadIdx.x; col < width; col += THREADS) {
        const float residual = y[col];
        output[col] =
            ((x[col] - mean) * inverse_std + residual) * residual;
    }
}

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(), "Expected matching devices");
    TORCH_CHECK(
        x.scalar_type() == at::kFloat && y.scalar_type() == at::kFloat,
        "Expected float32 tensors");
    TORCH_CHECK(
        x.dim() == 2 && y.sizes() == x.sizes(),
        "Expected matching two-dimensional tensors");
    TORCH_CHECK(x.size(1) > 1, "Expected more than one spatial element");

    const c10::cuda::CUDAGuard device_guard(x.device());
    x = x.contiguous();
    y = y.contiguous();

    auto output = torch::empty_like(x);
    if (x.size(0) == 0)
        return output;

    const auto stream = at::cuda::getCurrentCUDAStream();
    const dim3 grid(static_cast<unsigned int>(x.size(0)));

    const bool aligned =
        (reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(y.data_ptr<float>()) % 16 == 0);

    if (x.size(1) == 8192 && aligned) {
        fused_norm_8192<<<grid, THREADS, 0, stream>>>(
            x.data_ptr<float>(),
            y.data_ptr<float>(),
            output.data_ptr<float>(),
            static_cast<float>(eps));
    } else {
        fused_norm_generic<<<grid, THREADS, 0, stream>>>(
            x.data_ptr<float>(),
            y.data_ptr<float>(),
            output.data_ptr<float>(),
            x.size(1),
            static_cast<float>(eps));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="fused_row_instance_norm_residual_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):
        x = self.bmm(x)

        if (
            x.is_cuda
            and y.is_cuda
            and x.device == y.device
            and x.dtype == torch.float32
            and y.dtype == torch.float32
            and x.ndim == 2
            and x.shape == y.shape
            and x.shape[1] > 1
            and not (
                torch.is_grad_enabled() and (x.requires_grad or y.requires_grad)
            )
        ):
            # The original [N, 1, 1, F] normalization reduces over F.
            return _fused_ops.fused_norm_residual_cuda(
                x, y, self.instance_norm.eps
            )

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