```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>
#include <cstdint>
#include <climits>

__device__ __forceinline__ float block_sum(float value) {
    __shared__ float partial[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)
        partial[warp] = value;
    __syncthreads();

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

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

__device__ __forceinline__ float finish(
    float centered, float inv_std, float residual
) {
    const float normalized = __fmul_rn(centered, inv_std);
    return __fmul_rn(__fadd_rn(normalized, residual), residual);
}

template<int ITEMS>
__global__ void fused_norm_vector(
    const float4* __restrict__ input,
    const float4* __restrict__ residual,
    float4* __restrict__ output,
    int vectors,
    float epsilon
) {
    const int64_t base = int64_t(blockIdx.x) * vectors;
    float4 values[ITEMS];
    float sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        const float4 value = col < vectors
            ? input[base + col]
            : make_float4(0.0f, 0.0f, 0.0f, 0.0f);
        values[i] = value;
        sum += (value.x + value.y) + (value.z + value.w);
    }

    const float mean = block_sum(sum) / float(vectors * 4);
    float squared_sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < vectors) {
            float4 value = values[i];
            value.x = __fsub_rn(value.x, mean);
            value.y = __fsub_rn(value.y, mean);
            value.z = __fsub_rn(value.z, mean);
            value.w = __fsub_rn(value.w, mean);
            values[i] = value;

            squared_sum = fmaf(value.x, value.x, squared_sum);
            squared_sum = fmaf(value.y, value.y, squared_sum);
            squared_sum = fmaf(value.z, value.z, squared_sum);
            squared_sum = fmaf(value.w, value.w, squared_sum);
        }
    }

    const float inv_std = rsqrtf(
        block_sum(squared_sum) / float(vectors * 4) + epsilon);

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < vectors) {
            const float4 value = values[i];
            const float4 y = residual[base + col];
            output[base + col] = make_float4(
                finish(value.x, inv_std, y.x),
                finish(value.y, inv_std, y.y),
                finish(value.z, inv_std, y.z),
                finish(value.w, inv_std, y.w));
        }
    }
}

__global__ void fused_norm_scalar(
    const float* __restrict__ input,
    const float* __restrict__ residual,
    float* __restrict__ output,
    int columns,
    float epsilon
) {
    const int64_t base = int64_t(blockIdx.x) * columns;
    float values[32];
    float sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < 32; ++i) {
        const int col = threadIdx.x + i * 256;
        const float value = col < columns ? input[base + col] : 0.0f;
        values[i] = value;
        sum += value;
    }

    const float mean = block_sum(sum) / float(columns);
    float squared_sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < 32; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < columns) {
            values[i] = __fsub_rn(values[i], mean);
            squared_sum = fmaf(values[i], values[i], squared_sum);
        }
    }

    const float inv_std = rsqrtf(
        block_sum(squared_sum) / float(columns) + epsilon);

    #pragma unroll
    for (int i = 0; i < 32; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < columns) {
            output[base + col] = finish(
                values[i], inv_std, residual[base + col]);
        }
    }
}

torch::Tensor fused_norm_residual(
    torch::Tensor input,
    torch::Tensor residual,
    double epsilon
) {
    TORCH_CHECK(input.is_cuda() && residual.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.device() == residual.device(),
                "Devices must match");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                residual.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(input.is_contiguous() && residual.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(input.dim() == 2 && residual.dim() == 2 &&
                input.sizes() == residual.sizes(),
                "Expected matching matrices");
    TORCH_CHECK(input.size(0) <= INT_MAX &&
                input.size(1) > 1 && input.size(1) <= 8192,
                "Unsupported dimensions");

    const c10::cuda::CUDAGuard guard(input.device());
    const int rows = int(input.size(0));
    const int columns = int(input.size(1));
    auto output = torch::empty_like(input);

    if (rows == 0)
        return output;

    const auto stream = at::cuda::getCurrentCUDAStream();
    const float* z = input.data_ptr<float>();
    const float* y = residual.data_ptr<float>();
    float* out = output.data_ptr<float>();

    const bool aligned =
        ((reinterpret_cast<uintptr_t>(z) |
          reinterpret_cast<uintptr_t>(y) |
          reinterpret_cast<uintptr_t>(out)) & 15) == 0;

    if (columns % 4 == 0 && aligned) {
        const int vectors = columns / 4;
        const auto* zv = reinterpret_cast<const float4*>(z);
        const auto* yv = reinterpret_cast<const float4*>(y);
        auto* ov = reinterpret_cast<float4*>(out);

        if (vectors <= 256) {
            fused_norm_vector<1><<<rows, 256, 0, stream>>>(
                zv, yv, ov, vectors, float(epsilon));
        } else if (vectors <= 512) {
            fused_norm_vector<2><<<rows, 256, 0, stream>>>(
                zv, yv, ov, vectors, float(epsilon));
        } else if (vectors <= 1024) {
            fused_norm_vector<4><<<rows, 256, 0, stream>>>(
                zv, yv, ov, vectors, float(epsilon));
        } else {
            fused_norm_vector<8><<<rows, 256, 0, stream>>>(
                zv, yv, ov, vectors, float(epsilon));
        }
    } else {
        fused_norm_scalar<<<rows, 256, 0, stream>>>(
            z, y, out, columns, float(epsilon));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_ops = load_inline(
    name="linear_native_fused_instance_residual_v6",
    cpp_sources=r"""
torch::Tensor fused_norm_residual(
    torch::Tensor input,
    torch::Tensor residual,
    double epsilon);
""",
    cuda_sources=_cuda_source,
    functions=["fused_norm_residual"],
    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):
        # Preserve the reference GEMM precision and fuse its remaining operators.
        z = self.bmm(x)

        if (
            z.is_cuda
            and y.is_cuda
            and z.device == y.device
            and z.dtype == y.dtype == torch.float32
            and z.ndim == 2
            and y.shape == z.shape
            and z.shape[0] <= 2147483647
            and 1 < z.shape[1] <= 8192
        ):
            return _ops.fused_norm_residual(
                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
```