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

_cpp_source = """
#include <torch/extension.h>
torch::Tensor min_softmax_cuda(torch::Tensor x);
"""

_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>

__device__ __forceinline__ float minimum(float a, float b) {
    return (isnan(b) || b < a) ? b : a;
}

// Each warp processes 32 adjacent spatial positions.
// Eight warps cooperate across the 24 channels.
__global__ void min_softmax_kernel(
    const float* __restrict__ x,
    float* __restrict__ out,
    const int64_t spatial)
{
    constexpr int CHANNELS = 24;
    constexpr int DEPTH = 22;

    const int lane = threadIdx.x & 31;
    const int group = threadIdx.x >> 5;
    const int64_t pos = int64_t(blockIdx.x) * 32 + lane;
    const int64_t batch = blockIdx.y;
    const bool valid = pos < spatial;

    __shared__ float scratch[8][32];

    float m0 = CUDART_INF_F;
    float m1 = CUDART_INF_F;
    float m2 = CUDART_INF_F;

    if (valid) {
        const int64_t base =
            (batch * CHANNELS + group) * DEPTH * spatial + pos;
        const int64_t channel_step = 8LL * DEPTH * spatial;

        #pragma unroll
        for (int d = 0; d < DEPTH; ++d) {
            const int64_t offset = base + int64_t(d) * spatial;
            m0 = minimum(m0, x[offset]);
            m1 = minimum(m1, x[offset + channel_step]);
            m2 = minimum(m2, x[offset + 2 * channel_step]);
        }
    }

    scratch[group][lane] = fmaxf(m0, fmaxf(m1, m2));
    __syncthreads();

    float maximum = -CUDART_INF_F;
    #pragma unroll
    for (int g = 0; g < 8; ++g) {
        maximum = fmaxf(maximum, scratch[g][lane]);
    }

    // Complete all reads before reusing shared memory.
    __syncthreads();

    const float e0 = expf(m0 - maximum);
    const float e1 = expf(m1 - maximum);
    const float e2 = expf(m2 - maximum);

    scratch[group][lane] = e0 + e1 + e2;
    __syncthreads();

    float denominator = 0.0f;
    #pragma unroll
    for (int g = 0; g < 8; ++g) {
        denominator += scratch[g][lane];
    }

    if (valid) {
        const float inverse = 1.0f / denominator;
        const int64_t base =
            (batch * CHANNELS + group) * spatial + pos;

        out[base] = e0 * inverse;
        out[base + 8 * spatial] = e1 * inverse;
        out[base + 16 * spatial] = e2 * inverse;
    }
}

torch::Tensor min_softmax_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.dim() == 5 && x.size(1) == 24 && x.size(2) == 22,
                "Expected shape (N, 24, 22, H, W)");
    TORCH_CHECK(x.is_contiguous(), "Expected contiguous input");

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

    auto out = torch::empty(
        {x.size(0), x.size(1), x.size(3), x.size(4)}, x.options());

    const int64_t spatial = x.size(3) * x.size(4);
    if (out.numel() == 0) {
        return out;
    }

    const dim3 grid(
        static_cast<unsigned int>((spatial + 31) / 32),
        static_cast<unsigned int>(x.size(0)));

    min_softmax_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(x.data_ptr<float>(), out.data_ptr<float>(), spatial);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_min_softmax = load_inline(
    name="conv3d_depth_min_channel_softmax_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["min_softmax_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, dim):
        super().__init__()
        self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.dim = dim

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

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.is_contiguous()
            and self.dim in (2, -3)
            and x.shape[1] == 24
            and x.shape[2] == 22
            and x.shape[0] <= 65535
            and not (torch.is_grad_enabled() and x.requires_grad)
        ):
            return _min_softmax.min_softmax_cuda(x)

        return torch.softmax(torch.min(x, dim=self.dim)[0], dim=1)
```