```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(
    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.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cmath>

constexpr int STATS_THREADS = 256;
constexpr int STATS_TILE = 4096;

__device__ __forceinline__ void merge_stats(
    float& mean, float& m2, int& count,
    float other_mean, float other_m2, int other_count) {
    int total = count + other_count;
    if (total > 0) {
        float delta = other_mean - mean;
        float fraction = static_cast<float>(other_count) / total;
        m2 += other_m2 + delta * delta *
              (static_cast<float>(count) * other_count / total);
        mean += delta * fraction;
        count = total;
    }
}

__global__ void partial_stats_kernel(
    const float* __restrict__ x,
    float* __restrict__ partial,
    int group_size,
    int chunks) {
    int group = blockIdx.x;
    int chunk = blockIdx.y;
    int begin = chunk * STATS_TILE;
    int end = min(begin + STATS_TILE, group_size);

    float mean = 0.0f;
    float m2 = 0.0f;
    int count = 0;

    for (int i = begin + threadIdx.x; i < end; i += blockDim.x) {
        float value = x[static_cast<int64_t>(group) * group_size + i];
        ++count;
        float delta = value - mean;
        mean += delta / count;
        m2 += delta * (value - mean);
    }

    for (int offset = 16; offset > 0; offset >>= 1) {
        float other_mean = __shfl_down_sync(0xffffffff, mean, offset);
        float other_m2 = __shfl_down_sync(0xffffffff, m2, offset);
        int other_count = __shfl_down_sync(0xffffffff, count, offset);
        if ((threadIdx.x & 31) + offset < 32) {
            merge_stats(mean, m2, count, other_mean, other_m2, other_count);
        }
    }

    __shared__ float warp_mean[8];
    __shared__ float warp_m2[8];
    __shared__ int warp_count[8];

    int lane = threadIdx.x & 31;
    int warp = threadIdx.x >> 5;
    if (lane == 0) {
        warp_mean[warp] = mean;
        warp_m2[warp] = m2;
        warp_count[warp] = count;
    }
    __syncthreads();

    if (warp == 0) {
        mean = lane < 8 ? warp_mean[lane] : 0.0f;
        m2 = lane < 8 ? warp_m2[lane] : 0.0f;
        count = lane < 8 ? warp_count[lane] : 0;

        for (int offset = 16; offset > 0; offset >>= 1) {
            float other_mean = __shfl_down_sync(0xffffffff, mean, offset);
            float other_m2 = __shfl_down_sync(0xffffffff, m2, offset);
            int other_count = __shfl_down_sync(0xffffffff, count, offset);
            if (lane + offset < 32) {
                merge_stats(mean, m2, count, other_mean, other_m2, other_count);
            }
        }

        if (lane == 0) {
            int64_t index = (static_cast<int64_t>(group) * chunks + chunk) * 2;
            partial[index] = mean;
            partial[index + 1] = m2;
        }
    }
}

__global__ void finish_stats_kernel(
    const float* __restrict__ partial,
    float* __restrict__ stats,
    int total_groups,
    int group_size,
    int chunks,
    float eps) {
    int group = blockIdx.x * blockDim.x + threadIdx.x;
    if (group >= total_groups) return;

    float mean = 0.0f;
    float m2 = 0.0f;
    int count = 0;

    for (int chunk = 0; chunk < chunks; ++chunk) {
        int chunk_count = min(STATS_TILE, group_size - chunk * STATS_TILE);
        int64_t index = (static_cast<int64_t>(group) * chunks + chunk) * 2;
        merge_stats(
            mean, m2, count,
            partial[index], partial[index + 1], chunk_count);
    }

    stats[group * 2] = mean;
    stats[group * 2 + 1] = rsqrtf(fmaxf(m2 / group_size, 0.0f) + eps);
}

__global__ void fused_output_64_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    const float* __restrict__ stats,
    float* __restrict__ output,
    int spatial,
    int groups,
    int channels_per_group) {
    constexpr int PIXELS = 64;
    constexpr int CHANNEL_LANES = 4;

    int pixel_lane = threadIdx.x;
    int channel_lane = threadIdx.y;
    int pixel = blockIdx.x * PIXELS + pixel_lane;
    int batch = blockIdx.y;
    bool valid = pixel < spatial;

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

    #pragma unroll
    for (int k = 0; k < 16; ++k) {
        int channel = channel_lane + CHANNEL_LANES * k;
        float value = -CUDART_INF_F;
        if (valid) {
            int group = batch * groups + channel / channels_per_group;
            float conv = x[
                (static_cast<int64_t>(batch) * 64 + channel) * spatial + pixel];
            float normalized = (conv - stats[group * 2]) * stats[group * 2 + 1];
            normalized = normalized * weight[channel] + bias[channel];
            float t = tanhf(normalized);
            value = conv + t * (t + 3.0f) * (1.0f / 6.0f);
        }
        values[k] = value;
        local_max = fmaxf(local_max, value);
    }

    __shared__ float scratch[CHANNEL_LANES][PIXELS];
    scratch[channel_lane][pixel_lane] = local_max;
    __syncthreads();

    float maximum = scratch[0][pixel_lane];
    #pragma unroll
    for (int lane = 1; lane < CHANNEL_LANES; ++lane) {
        maximum = fmaxf(maximum, scratch[lane][pixel_lane]);
    }
    __syncthreads();

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

    scratch[channel_lane][pixel_lane] = sum;
    __syncthreads();

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

torch::Tensor fused_post_conv(
    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.is_contiguous() && x.dim() == 4 && x.size(1) == 64,
                "Expected contiguous NCHW input with 64 channels");
    TORCH_CHECK(groups > 0 && 64 % groups == 0, "Invalid group count");

    c10::cuda::CUDAGuard guard(x.device());
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    int batch = static_cast<int>(x.size(0));
    int spatial = static_cast<int>(x.size(2) * x.size(3));
    int channels_per_group = 64 / groups;
    int group_size = channels_per_group * spatial;
    int total_groups = batch * groups;
    int chunks = (group_size + STATS_TILE - 1) / STATS_TILE;

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

    auto partial = torch::empty({total_groups, chunks, 2}, x.options());
    auto stats = torch::empty({total_groups, 2}, x.options());

    partial_stats_kernel<<<dim3(total_groups, chunks), STATS_THREADS, 0, stream>>>(
        x.data_ptr<float>(), partial.data_ptr<float>(), group_size, chunks);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    finish_stats_kernel<<<(total_groups + 255) / 256, 256, 0, stream>>>(
        partial.data_ptr<float>(), stats.data_ptr<float>(),
        total_groups, group_size, chunks, static_cast<float>(eps));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    fused_output_64_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), channels_per_group);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_fused_ops = load_inline(
    name="groupnorm_tanh_hardswish_residual_lse_cuda_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_post_conv"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    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 self.group_norm.weight.dtype == torch.float32
            and self.group_norm.bias.dtype == torch.float32
        ):
            return _fused_ops.fused_post_conv(
                x_conv.contiguous(),
                self.group_norm.weight,
                self.group_norm.bias,
                self.group_norm.num_groups,
                self.group_norm.eps,
            )

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