```python
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor fused_post_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t groups,
    double eps);
"""

_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 <math_constants.h>

constexpr int TILE = 4096;

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

__global__ void partial_moments_kernel(
    const float* __restrict__ x,
    float* __restrict__ partial,
    int group_size,
    int parts)
{
    const int group = blockIdx.x;
    const int part = blockIdx.y;
    const int tid = threadIdx.x;
    const int start = part * TILE;
    const int end = min(start + TILE, group_size);
    const int64_t base = static_cast<int64_t>(group) * group_size;

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

    for (int i = start + tid; i < end; i += blockDim.x) {
        const float v = x[base + i];
        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];

    const int lane = tid & 31;
    const int warp = tid >> 5;

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

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

        if (lane == 0) {
            const int64_t offset =
                (static_cast<int64_t>(group) * parts + part) * 2;
            partial[offset] = sum;
            partial[offset + 1] = square_sum;
        }
    }
}

__global__ void finish_moments_kernel(
    const float* __restrict__ partial,
    float* __restrict__ stats,
    int group_size,
    int parts,
    float eps)
{
    const int group = blockIdx.x;
    const int lane = threadIdx.x;

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

    for (int part = lane; part < parts; part += 32) {
        const int64_t offset =
            (static_cast<int64_t>(group) * parts + part) * 2;
        sum += partial[offset];
        square_sum += partial[offset + 1];
    }

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

    if (lane == 0) {
        const float mean = sum / static_cast<float>(group_size);
        const float variance = fmaxf(
            fmaf(-mean, mean, square_sum / static_cast<float>(group_size)),
            0.0f);
        stats[group * 2] = mean;
        stats[group * 2 + 1] = rsqrtf(variance + eps);
    }
}

__global__ void fused_activation_logsumexp_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    const float* __restrict__ stats,
    float* __restrict__ output,
    int spatial,
    int groups)
{
    // Four channel partitions cooperate on 64 spatial positions.
    const int pixel_lane = threadIdx.x;
    const int channel_partition = threadIdx.y;
    const int pixel = blockIdx.x * 64 + pixel_lane;
    const int batch = blockIdx.y;
    const bool valid = pixel < spatial;
    const int channels_per_group = 64 / groups;

    __shared__ float maxima[4][64];
    __shared__ float sums[4][64];

    float values[16];
    float local_max = -CUDART_INF_F;

    #pragma unroll
    for (int i = 0; i < 16; ++i) {
        float value = -CUDART_INF_F;
        if (valid) {
            const int channel = channel_partition * 16 + i;
            const int group = batch * groups + channel / channels_per_group;
            const float mean = stats[group * 2];
            const float inverse_std = stats[group * 2 + 1];
            const int64_t index =
                (static_cast<int64_t>(batch) * 64 + channel) * spatial + pixel;

            const float residual = x[index];
            const float normalized =
                (residual - mean) * inverse_std;
            const float affine =
                fmaf(normalized, weight[channel], bias[channel]);
            const float t = tanhf(affine);

            // Tanh is in [-1, 1], so Hardswish's clamp is inactive.
            const float activated = (t * (t + 3.0f)) * (1.0f / 6.0f);
            value = residual + activated;
        }
        values[i] = value;
        local_max = fmaxf(local_max, value);
    }

    maxima[channel_partition][pixel_lane] = local_max;
    __syncthreads();

    float maximum = maxima[0][pixel_lane];
    #pragma unroll
    for (int partition = 1; partition < 4; ++partition)
        maximum = fmaxf(maximum, maxima[partition][pixel_lane]);

    float sum = 0.0f;
    if (valid) {
        #pragma unroll
        for (int i = 0; i < 16; ++i)
            sum += __expf(values[i] - maximum);
    }

    sums[channel_partition][pixel_lane] = sum;
    __syncthreads();

    if (channel_partition == 0 && valid) {
        const float total =
            (sums[0][pixel_lane] + sums[1][pixel_lane]) +
            (sums[2][pixel_lane] + sums[3][pixel_lane]);
        output[static_cast<int64_t>(batch) * spatial + pixel] =
            maximum + logf(total);
    }
}

torch::Tensor fused_post_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t groups,
    double eps)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 tensor");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == 64 && x.is_contiguous(),
                "Expected contiguous NCHW input with 64 channels");
    TORCH_CHECK(groups > 0 && 64 % groups == 0, "Invalid group count");
    TORCH_CHECK(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,
                "Invalid normalization parameters");

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t spatial64 = x.size(2) * x.size(3);
    TORCH_CHECK(spatial64 > 0 &&
                spatial64 * (64 / groups) <= 2147483647LL &&
                x.size(0) <= 65535,
                "Input dimensions exceed kernel limits");

    const int batch = static_cast<int>(x.size(0));
    const int spatial = static_cast<int>(spatial64);
    const int group_size = (64 / static_cast<int>(groups)) * spatial;
    const int parts = (group_size - 1) / TILE + 1;
    const int total_groups = batch * static_cast<int>(groups);

    auto output = torch::empty({batch, 1, x.size(2), x.size(3)}, x.options());
    if (batch == 0)
        return output;

    TORCH_CHECK(parts <= 65535, "Too many reduction tiles");

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

    partial_moments_kernel<<<dim3(total_groups, parts), 256, 0, stream>>>(
        x.data_ptr<float>(), partial.data_ptr<float>(), group_size, parts);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    finish_moments_kernel<<<total_groups, 32, 0, stream>>>(
        partial.data_ptr<float>(), stats.data_ptr<float>(),
        group_size, parts, static_cast<float>(eps));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    fused_activation_logsumexp_kernel<<<
        dim3((spatial + 63) / 64, batch), dim3(64, 4), 0, stream>>>(
        x.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
        stats.data_ptr<float>(), output.data_ptr<float>(),
        spatial, static_cast<int>(groups));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_fused_ops = load_inline(
    name="fused_gn_tanh_hardswish_residual_lse_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_post_conv_cuda"],
    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.shape[1] == 64
            and x_conv.is_contiguous()
            and self.group_norm.weight.dtype == torch.float32
            and self.group_norm.bias.dtype == torch.float32
        ):
            return _fused_ops.fused_post_conv_cuda(
                x_conv,
                self.group_norm.weight,
                self.group_norm.bias,
                self.group_norm.num_groups,
                self.group_norm.eps,
            )

        x_norm = self.group_norm(x_conv)
        activated = F.hardswish(torch.tanh(x_norm))
        return torch.logsumexp(x_conv + activated, dim=1, keepdim=True)
```