```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 <algorithm>
#include <vector>

constexpr int THREADS = 256;
constexpr int CHUNK = 2048;

__device__ __forceinline__ float warp_sum(float v) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xffffffff, v, offset);
    return v;
}

// Each block computes statistics for a contiguous portion of one group.
// Centering around a group element reduces cancellation in the variance.
__global__ void partial_stats(
    const float* __restrict__ x,
    float2* __restrict__ partial,
    int group_elements,
    int chunks)
{
    const int group = blockIdx.x;
    const int chunk = blockIdx.y;
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;
    const long long base = (long long)group * group_elements;
    const float center = x[base];

    float sum = 0.0f;
    float square_sum = 0.0f;

    #pragma unroll
    for (int j = 0; j < CHUNK / THREADS; ++j) {
        const int i = chunk * CHUNK + threadIdx.x + j * THREADS;
        if (i < group_elements) {
            const float v = x[base + i] - center;
            sum += v;
            square_sum = fmaf(v, v, square_sum);
        }
    }

    sum = warp_sum(sum);
    square_sum = warp_sum(square_sum);

    __shared__ float sums[8];
    __shared__ float squares[8];

    if (lane == 0) {
        sums[warp] = sum;
        squares[warp] = square_sum;
    }
    __syncthreads();

    if (warp == 0) {
        sum = lane < 8 ? sums[lane] : 0.0f;
        square_sum = lane < 8 ? squares[lane] : 0.0f;
        sum = warp_sum(sum);
        square_sum = warp_sum(square_sum);

        if (lane == 0)
            partial[(long long)group * chunks + chunk] =
                make_float2(sum, square_sum);
    }
}

__global__ void finish_stats(
    const float* __restrict__ x,
    const float2* __restrict__ partial,
    float2* __restrict__ stats,
    int group_elements,
    int chunks,
    float eps)
{
    const int group = blockIdx.x;
    float sum = 0.0f;
    float square_sum = 0.0f;

    for (int i = threadIdx.x; i < chunks; i += 32) {
        const float2 p = partial[(long long)group * chunks + i];
        sum += p.x;
        square_sum += p.y;
    }

    sum = warp_sum(sum);
    square_sum = warp_sum(square_sum);

    if (threadIdx.x == 0) {
        const float shifted_mean = sum / (float)group_elements;
        const float variance = fmaxf(
            fmaf(-shifted_mean, shifted_mean,
                 square_sum / (float)group_elements),
            0.0f);
        const float mean =
            x[(long long)group * group_elements] + shifted_mean;
        stats[group] = make_float2(mean, rsqrtf(variance + eps));
    }
}

// Neighboring threads handle neighboring spatial positions, so all
// channel reads are coalesced. Intermediate activations stay in registers.
__global__ void fused_epilogue(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    const float2* __restrict__ stats,
    float* __restrict__ output,
    int spatial,
    long long output_elements)
{
    const long long index =
        (long long)blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= output_elements)
        return;

    const int batch = index / spatial;
    const int position = index % spatial;
    const long long base = (long long)batch * 64 * spatial + position;

    float values[64];
    float maximum = -CUDART_INF_F;

    #pragma unroll
    for (int group = 0; group < 16; ++group) {
        const float2 s = stats[batch * 16 + group];

        #pragma unroll
        for (int k = 0; k < 4; ++k) {
            const int channel = group * 4 + k;
            const float residual = x[base + (long long)channel * spatial];
            const float normalized = (residual - s.x) * s.y;
            const float affine =
                normalized * weight[channel] + bias[channel];
            const float t = tanhf(affine);

            // tanh lies in [-1, 1], so the Hardswish clamp is inactive.
            const float activation = (t * (t + 3.0f)) / 6.0f;
            const float value = residual + activation;
            values[channel] = value;
            maximum = fmaxf(maximum, value);
        }
    }

    float sum = 0.0f;
    #pragma unroll
    for (int channel = 0; channel < 64; ++channel)
        sum += expf(values[channel] - maximum);

    output[index] = maximum + logf(sum);
}

torch::Tensor fused_post_conv(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    double eps)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected CUDA float32 input");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == 64 && x.is_contiguous(),
                "Expected contiguous NCHW input with 64 channels");
    TORCH_CHECK(weight.is_cuda() && bias.is_cuda() &&
                weight.device() == x.device() &&
                bias.device() == x.device() &&
                weight.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat &&
                weight.is_contiguous() && bias.is_contiguous() &&
                weight.numel() == 64 && bias.numel() == 64,
                "Expected contiguous float32 affine parameters");

    const c10::cuda::CUDAGuard guard(x.device());
    const int batch = x.size(0);
    const int height = x.size(2);
    const int width = x.size(3);
    const int spatial = height * width;
    const int group_elements = 4 * spatial;
    const int chunks = (group_elements + CHUNK - 1) / CHUNK;
    const int total_groups = batch * 16;

    auto output = torch::empty({batch, 1, height, width}, x.options());
    if (output.numel() == 0)
        return output;

    auto partial = torch::empty({total_groups, chunks, 2}, x.options());
    auto stats = torch::empty({total_groups, 2}, x.options());
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    partial_stats<<<dim3(total_groups, chunks), THREADS, 0, stream>>>(
        x.data_ptr<float>(),
        reinterpret_cast<float2*>(partial.data_ptr<float>()),
        group_elements,
        chunks);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    finish_stats<<<total_groups, 32, 0, stream>>>(
        x.data_ptr<float>(),
        reinterpret_cast<const float2*>(partial.data_ptr<float>()),
        reinterpret_cast<float2*>(stats.data_ptr<float>()),
        group_elements,
        chunks,
        static_cast<float>(eps));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    const long long count = (long long)batch * spatial;
    fused_epilogue<<<(count + 127) / 128, 128, 0, stream>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias.data_ptr<float>(),
        reinterpret_cast<const float2*>(stats.data_ptr<float>()),
        output.data_ptr<float>(),
        spatial,
        count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_extension = load_inline(
    name="fused_groupnorm_hardswish_residual_lse_v1",
    cpp_sources="""
#include <torch/extension.h>
torch::Tensor fused_post_conv(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    double eps);
""",
    cuda_sources=_cuda_source,
    functions=["fused_post_conv"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self, in_channels, out_channels, kernel_size, groups, eps=1e-5
    ):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.group_norm = nn.GroupNorm(groups, out_channels, eps=eps)
        self.tanh = nn.Tanh()
        self.hard_swish = nn.Hardswish()

    def forward(self, x):
        x_conv = self.conv(x)

        if (
            x_conv.is_cuda
            and x_conv.dtype == torch.float32
            and x_conv.is_contiguous()
            and x_conv.shape[1] == 64
            and self.group_norm.num_groups == 16
            and self.group_norm.weight.dtype == torch.float32
            and self.group_norm.bias.dtype == torch.float32
        ):
            return _extension.fused_post_conv(
                x_conv,
                self.group_norm.weight,
                self.group_norm.bias,
                self.group_norm.eps,
            )

        x_norm = self.group_norm(x_conv)
        x_activated = self.hard_swish(self.tanh(x_norm))
        return torch.logsumexp(x_conv + x_activated, dim=1, keepdim=True)
```