```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 <ATen/cuda/CUDABlas.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_fp16.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <climits>

__global__ void split_inputs(
    const float* __restrict__ weight,
    const float* __restrict__ x,
    __half* __restrict__ packed,
    int64_t weight_count,
    int64_t total
) {
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < total;
         i += int64_t(blockDim.x) * gridDim.x) {
        const float v = i < weight_count ? weight[i] : x[i - weight_count];
        const __half hi = __float2half_rn(v);
        const float remainder = (v - __half2float(hi)) * 2048.0f;
        packed[i] = hi;
        packed[total + i] = __float2half_rn(remainder);
    }
}

__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; 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; offset >>= 1)
            value += __shfl_down_sync(0xffffffff, value, offset);
        if (lane == 0)
            partial[0] = value;
    }
    __syncthreads();
    const float result = partial[0];
    __syncthreads();
    return result;
}

template<int ITEMS>
__global__ void fused_epilogue(
    const float* __restrict__ linear,
    const float* __restrict__ bias,
    const float* __restrict__ residual,
    float* __restrict__ output,
    int columns,
    float epsilon
) {
    const int64_t base = int64_t(blockIdx.x) * columns;
    float values[ITEMS];
    float sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        const float value = col < columns
            ? __fadd_rn(linear[base + col], bias[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 < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < columns) {
            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 < ITEMS; ++i) {
        const int col = threadIdx.x + i * 256;
        if (col < columns) {
            const float y = residual[base + col];
            const float normalized = __fmul_rn(values[i], inv_std);
            output[base + col] =
                __fmul_rn(__fadd_rn(normalized, y), y);
        }
    }
}

struct MathModeGuard {
    cublasHandle_t handle;
    cublasMath_t previous;

    explicit MathModeGuard(cublasHandle_t h) : handle(h) {
        TORCH_CHECK(
            cublasGetMathMode(handle, &previous) == CUBLAS_STATUS_SUCCESS,
            "Cannot query cuBLAS math mode");
        const auto mode = static_cast<cublasMath_t>(
            static_cast<unsigned int>(previous) |
            static_cast<unsigned int>(
                CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION));
        TORCH_CHECK(
            cublasSetMathMode(handle, mode) == CUBLAS_STATUS_SUCCESS,
            "Cannot set cuBLAS math mode");
    }

    ~MathModeGuard() {
        cublasSetMathMode(handle, previous);
    }
};

torch::Tensor split_linear_norm(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    torch::Tensor y,
    double epsilon
) {
    TORCH_CHECK(
        x.is_cuda() && weight.is_cuda() && bias.is_cuda() && y.is_cuda(),
        "Expected CUDA tensors");
    TORCH_CHECK(
        x.device() == weight.device() &&
        x.device() == bias.device() && x.device() == y.device(),
        "Devices must match");
    TORCH_CHECK(
        x.scalar_type() == at::kFloat &&
        weight.scalar_type() == at::kFloat &&
        bias.scalar_type() == at::kFloat &&
        y.scalar_type() == at::kFloat,
        "Expected float32 tensors");
    TORCH_CHECK(
        x.is_contiguous() && weight.is_contiguous() &&
        bias.is_contiguous() && y.is_contiguous(),
        "Expected contiguous tensors");
    TORCH_CHECK(
        x.dim() == 2 && weight.dim() == 2 &&
        bias.dim() == 1 && y.dim() == 2,
        "Invalid tensor ranks");
    TORCH_CHECK(
        x.size(1) == weight.size(1) &&
        bias.size(0) == weight.size(0) &&
        y.size(0) == x.size(0) && y.size(1) == weight.size(0),
        "Incompatible shapes");
    TORCH_CHECK(
        x.size(0) <= INT_MAX &&
        x.size(1) > 0 && x.size(1) <= INT_MAX &&
        weight.size(0) > 1 && weight.size(0) <= 8192,
        "Unsupported dimensions");

    const c10::cuda::CUDAGuard device_guard(x.device());
    const int rows = int(x.size(0));
    const int columns = int(weight.size(0));
    const int inner = int(x.size(1));

    auto output = torch::empty({rows, columns}, x.options());
    if (rows == 0)
        return output;

    auto linear = torch::empty_like(output);
    const int64_t weight_count = weight.numel();
    const int64_t total = weight_count + x.numel();
    auto packed = torch::empty(
        {2, total}, x.options().dtype(at::kHalf));

    auto* data = reinterpret_cast<__half*>(
        packed.data_ptr<at::Half>());
    const auto stream = at::cuda::getCurrentCUDAStream();
    const int64_t requested_blocks = (total + 255) / 256;
    const int blocks = int(requested_blocks < 65535
        ? requested_blocks : 65535);

    split_inputs<<<blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(), x.data_ptr<float>(),
        data, weight_count, total);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    const __half* weight_hi = data;
    const __half* x_hi = data + weight_count;
    const __half* weight_lo = data + total;
    const __half* x_lo = data + total + weight_count;

    cublasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();

    {
        MathModeGuard math_guard(handle);

        const float one = 1.0f;
        const float zero = 0.0f;
        const float correction_scale = 1.0f / 2048.0f;

        // Two half components preserve substantially more input precision
        // than a single TF32 conversion. All GEMMs accumulate in float32.
        auto gemm = [&](const __half* w, const __half* a,
                        const float* alpha, const float* beta) {
            const cublasStatus_t status = cublasGemmEx(
                handle,
                CUBLAS_OP_T, CUBLAS_OP_N,
                columns, rows, inner,
                alpha,
                w, CUDA_R_16F, inner,
                a, CUDA_R_16F, inner,
                beta,
                linear.data_ptr<float>(), CUDA_R_32F, columns,
                CUBLAS_COMPUTE_32F,
                CUBLAS_GEMM_DEFAULT_TENSOR_OP);
            TORCH_CHECK(
                status == CUBLAS_STATUS_SUCCESS,
                "cublasGemmEx failed: ", int(status));
        };

        gemm(weight_lo, x_hi, &correction_scale, &zero);
        gemm(weight_hi, x_lo, &correction_scale, &one);
        gemm(weight_hi, x_hi, &one, &one);
    }

    if (columns <= 2048) {
        fused_epilogue<8><<<rows, 256, 0, stream>>>(
            linear.data_ptr<float>(), bias.data_ptr<float>(),
            y.data_ptr<float>(), output.data_ptr<float>(),
            columns, float(epsilon));
    } else if (columns <= 4096) {
        fused_epilogue<16><<<rows, 256, 0, stream>>>(
            linear.data_ptr<float>(), bias.data_ptr<float>(),
            y.data_ptr<float>(), output.data_ptr<float>(),
            columns, float(epsilon));
    } else {
        fused_epilogue<32><<<rows, 256, 0, stream>>>(
            linear.data_ptr<float>(), bias.data_ptr<float>(),
            y.data_ptr<float>(), output.data_ptr<float>(),
            columns, float(epsilon));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_ops = load_inline(
    name="split_half_linear_fused_norm_v3",
    cpp_sources=r"""
torch::Tensor split_linear_norm(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    torch::Tensor y,
    double epsilon);
""",
    cuda_sources=_cuda_source,
    functions=["split_linear_norm"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=["-lcublas"],
    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):
        weight = self.bmm.weight
        bias = self.bmm.bias

        if (
            x.is_cuda
            and y.is_cuda
            and x.device == y.device == weight.device == bias.device
            and x.dtype == y.dtype == weight.dtype == bias.dtype == torch.float32
            and x.ndim == 2
            and y.ndim == 2
            and x.shape[1] == weight.shape[1]
            and x.shape[1] > 0
            and y.shape == (x.shape[0], weight.shape[0])
            and 1 < weight.shape[0] <= 8192
            and not torch.is_autocast_enabled()
        ):
            return _ops.split_linear_norm(
                x.contiguous(),
                weight.contiguous(),
                bias.contiguous(),
                y.contiguous(),
                self.instance_norm.eps,
            )

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