```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 norm_residual_mul_cuda(
    torch::Tensor x, torch::Tensor y, double eps);
"""

_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.h>
#include <cuda_runtime.h>

template <typename T>
__device__ __forceinline__ T warp_sum(T value) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        value += __shfl_down_sync(0xffffffff, value, offset);
    }
    return value;
}

template <typename T>
__device__ __forceinline__ T block_sum(T value, T* shared) {
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;

    value = warp_sum(value);
    if (lane == 0) {
        shared[warp] = value;
    }
    __syncthreads();

    if (warp == 0) {
        value = lane < 8 ? shared[lane] : T(0);
        value = warp_sum(value);
        if (lane == 0) {
            shared[0] = value;
        }
    }
    __syncthreads();

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

template <typename T>
__device__ __forceinline__ T inverse_sqrt(T value);

template <>
__device__ __forceinline__ float inverse_sqrt<float>(float value) {
    return rsqrtf(value);
}

template <>
__device__ __forceinline__ double inverse_sqrt<double>(double value) {
    return rsqrt(value);
}

// CACHE_ITEMS=32 specializes the supplied 8192-element rows.
// CACHE_ITEMS=0 supports arbitrary output widths.
template <typename scalar_t, int CACHE_ITEMS>
__global__ void norm_residual_mul_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ y,
    scalar_t* __restrict__ output,
    int64_t width,
    double eps
) {
    using acc_t = at::acc_type<scalar_t, true>;

    __shared__ acc_t shared[8];
    acc_t cached[CACHE_ITEMS > 0 ? CACHE_ITEMS : 1];

    const int64_t base = static_cast<int64_t>(blockIdx.x) * width;
    acc_t sum = acc_t(0);

    if constexpr (CACHE_ITEMS > 0) {
        #pragma unroll
        for (int i = 0; i < CACHE_ITEMS; ++i) {
            const int column = threadIdx.x + i * 256;
            cached[i] = static_cast<acc_t>(x[base + column]);
            sum += cached[i];
        }
    } else {
        for (int64_t column = threadIdx.x; column < width; column += 256) {
            sum += static_cast<acc_t>(x[base + column]);
        }
    }

    const acc_t mean = block_sum(sum, shared) / static_cast<acc_t>(width);
    acc_t squared_deviations = acc_t(0);

    if constexpr (CACHE_ITEMS > 0) {
        #pragma unroll
        for (int i = 0; i < CACHE_ITEMS; ++i) {
            const acc_t difference = cached[i] - mean;
            squared_deviations += difference * difference;
        }
    } else {
        for (int64_t column = threadIdx.x; column < width; column += 256) {
            const acc_t difference = static_cast<acc_t>(x[base + column]) - mean;
            squared_deviations += difference * difference;
        }
    }

    const acc_t variance =
        block_sum(squared_deviations, shared) / static_cast<acc_t>(width);
    const acc_t inv_std = inverse_sqrt(variance + static_cast<acc_t>(eps));

    if constexpr (CACHE_ITEMS > 0) {
        #pragma unroll
        for (int i = 0; i < CACHE_ITEMS; ++i) {
            const int64_t index = base + threadIdx.x + i * 256;
            const acc_t residual = static_cast<acc_t>(y[index]);
            const scalar_t normalized =
                static_cast<scalar_t>((cached[i] - mean) * inv_std);
            const scalar_t added =
                static_cast<scalar_t>(static_cast<acc_t>(normalized) + residual);
            output[index] =
                static_cast<scalar_t>(static_cast<acc_t>(added) * residual);
        }
    } else {
        for (int64_t column = threadIdx.x; column < width; column += 256) {
            const int64_t index = base + column;
            const acc_t residual = static_cast<acc_t>(y[index]);
            const scalar_t normalized = static_cast<scalar_t>(
                (static_cast<acc_t>(x[index]) - mean) * inv_std);
            const scalar_t added =
                static_cast<scalar_t>(static_cast<acc_t>(normalized) + residual);
            output[index] =
                static_cast<scalar_t>(static_cast<acc_t>(added) * residual);
        }
    }
}

torch::Tensor norm_residual_mul_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(), "Tensor devices must match");
    TORCH_CHECK(x.scalar_type() == y.scalar_type(), "Tensor dtypes must match");
    TORCH_CHECK(x.dim() == 2 && x.sizes() == y.sizes(),
                "Expected matching two-dimensional tensors");
    TORCH_CHECK(x.is_contiguous() && y.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.size(1) > 1, "Normalization requires more than one element");

    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(x.get_device());
    const int64_t width = x.size(1);

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        x.scalar_type(),
        "norm_residual_mul_cuda",
        [&] {
            if (width == 8192) {
                norm_residual_mul_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>(),
                        width,
                        eps);
            } else {
                norm_residual_mul_kernel<scalar_t, 0>
                    <<<x.size(0), 256, 0, stream>>>(
                        x.data_ptr<scalar_t>(),
                        y.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(),
                        width,
                        eps);
            }
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="linear_instance_norm_residual_mul_ext",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["norm_residual_mul_cuda"],
    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 y.is_cuda
            and x.device == y.device
            and x.dtype == y.dtype
            and x.ndim == 2
            and x.shape == y.shape
        ):
            return _fused_ops.norm_residual_mul_cuda(
                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
```