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

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

torch::Tensor depth_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.h>
#include <cuda_runtime.h>
#include <stdint.h>

template<int CHANNELS, int THREADS>
__global__ void depth_min_softmax_kernel(
    const float* __restrict__ x,
    float* __restrict__ output,
    int depth,
    int64_t spatial,
    int64_t tiles_per_batch
) {
    const int64_t batch = blockIdx.x / tiles_per_batch;
    const int64_t tile = blockIdx.x % tiles_per_batch;
    const int64_t position = tile * THREADS + threadIdx.x;

    if (position >= spatial) {
        return;
    }

    const int64_t channel_stride = static_cast<int64_t>(depth) * spatial;
    const float* input = x + batch * CHANNELS * channel_stride + position;

    float minima[CHANNELS];

    #pragma unroll
    for (int c = 0; c < CHANNELS; ++c) {
        minima[c] = CUDART_INF_F;
    }

    #pragma unroll 1
    for (int d = 0; d < depth; ++d) {
        #pragma unroll
        for (int c = 0; c < CHANNELS; ++c) {
            const float value = input[c * channel_stride + d * spatial];
            // Preserve NaNs, matching torch.min.
            minima[c] =
                (value < minima[c] || isnan(value)) ? value : minima[c];
        }
    }

    float maximum = -CUDART_INF_F;

    #pragma unroll
    for (int c = 0; c < CHANNELS; ++c) {
        maximum = fmaxf(maximum, minima[c]);
    }

    float denominator = 0.0f;

    #pragma unroll
    for (int c = 0; c < CHANNELS; ++c) {
        minima[c] = expf(minima[c] - maximum);
        denominator += minima[c];
    }

    const float reciprocal = 1.0f / denominator;
    float* destination = output + batch * CHANNELS * spatial + position;

    #pragma unroll
    for (int c = 0; c < CHANNELS; ++c) {
        destination[c * spatial] = minima[c] * reciprocal;
    }
}

torch::Tensor depth_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, "Expected an NCDHW tensor");
    TORCH_CHECK(x.is_contiguous(), "Expected contiguous input");
    TORCH_CHECK(x.size(1) == 24, "Expected 24 channels");
    TORCH_CHECK(x.size(2) > 0, "Depth must be positive");

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

    const int64_t batch = x.size(0);
    const int depth = static_cast<int>(x.size(2));
    const int64_t height = x.size(3);
    const int64_t width = x.size(4);
    const int64_t spatial = height * width;

    auto output = torch::empty({batch, 24, height, width}, x.options());

    if (output.numel() == 0) {
        return output;
    }

    constexpr int threads = 128;
    const int64_t tiles = (spatial + threads - 1) / threads;
    const int64_t blocks = batch * tiles;

    TORCH_CHECK(blocks <= 2147483647LL, "Input is too large");

    depth_min_softmax_kernel<24, threads>
        <<<static_cast<unsigned int>(blocks), threads, 0,
           at::cuda::getCurrentCUDAStream()>>>(
            x.data_ptr<float>(),
            output.data_ptr<float>(),
            depth,
            spatial,
            tiles
        );

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="conv3d_depth_min_channel_softmax_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depth_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.size(1) == 24
            and self.dim in (2, -3)
            and x.is_contiguous()
        ):
            return _fused_ops.depth_min_softmax_cuda(x)

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