```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/AccumulateType.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

template <typename T>
__device__ __forceinline__ T block_sum(T value) {
    __shared__ T warp_sums[8];

    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)
        warp_sums[warp] = value;
    __syncthreads();

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

    T result = warp_sums[0];
    __syncthreads();
    return result;
}

template <typename scalar_t, int ITEMS>
__global__ void fused_norm_residual_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ y,
    scalar_t* __restrict__ output,
    int64_t columns,
    double epsilon
) {
    using acc_t = at::acc_type<scalar_t, true>;
    const int64_t base = int64_t(blockIdx.x) * columns;
    const int tid = threadIdx.x;

    acc_t values[ITEMS];
    acc_t sum = acc_t(0);

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int64_t column = tid + int64_t(i) * blockDim.x;
        values[i] = column < columns ? acc_t(x[base + column]) : acc_t(0);
        sum += values[i];
    }

    const acc_t mean = block_sum(sum) / acc_t(columns);
    acc_t squared_sum = acc_t(0);

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int64_t column = tid + int64_t(i) * blockDim.x;
        if (column < columns) {
            values[i] -= mean;
            squared_sum += values[i] * values[i];
        }
    }

    const acc_t variance = block_sum(squared_sum) / acc_t(columns);
    const acc_t inverse_std = acc_t(1) / sqrt(variance + acc_t(epsilon));

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int64_t column = tid + int64_t(i) * blockDim.x;
        if (column < columns) {
            const acc_t residual = acc_t(y[base + column]);
            const scalar_t normalized = scalar_t(values[i] * inverse_std);
            const scalar_t added = scalar_t(acc_t(normalized) + residual);
            output[base + column] = scalar_t(acc_t(added) * residual);
        }
    }
}

torch::Tensor fused_norm_residual(
    torch::Tensor x,
    torch::Tensor y,
    double epsilon
) {
    TORCH_CHECK(x.is_cuda() && y.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.device() == y.device(), "Devices must match");
    TORCH_CHECK(x.scalar_type() == y.scalar_type(), "Dtypes must match");
    TORCH_CHECK(x.dim() == 2 && x.sizes() == y.sizes(), "Expected matching matrices");
    TORCH_CHECK(x.is_contiguous() && y.is_contiguous(), "Expected contiguous tensors");
    TORCH_CHECK(x.size(1) > 1 && x.size(1) <= 8192, "Unsupported row width");

    const c10::cuda::CUDAGuard device_guard(x.device());
    auto output = torch::empty_like(x);
    if (x.size(0) == 0)
        return output;

    const auto stream = at::cuda::getCurrentCUDAStream();
    const int64_t columns = x.size(1);

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        x.scalar_type(), "fused_norm_residual", [&] {
            if (columns <= 2048) {
                fused_norm_residual_kernel<scalar_t, 8>
                    <<<x.size(0), 256, 0, stream>>>(
                        x.data_ptr<scalar_t>(), y.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(), columns, epsilon);
            } else {
                fused_norm_residual_kernel<scalar_t, 32>
                    <<<x.size(0), 256, 0, stream>>>(
                        x.data_ptr<scalar_t>(), y.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(), columns, epsilon);
            }
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="fused_instance_norm_residual_mul_v1",
    cpp_sources=[
        "torch::Tensor fused_norm_residual("
        "torch::Tensor x, torch::Tensor y, double epsilon);"
    ],
    cuda_sources=_cuda_source,
    functions=["fused_norm_residual"],
    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 == y.dtype
            and x.dtype in (
                torch.float16, torch.bfloat16, torch.float32, torch.float64
            )
            and x.ndim == 2
            and x.shape == y.shape
            and 1 < x.shape[1] <= 8192
        ):
            return _fused_ops.fused_norm_residual(
                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
```