```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 min_softmax_cuda(torch::Tensor x);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/AccumulateType.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>

__device__ __forceinline__ float exponential(float x) {
    return expf(x);
}

__device__ __forceinline__ double exponential(double x) {
    return exp(x);
}

template <typename T>
__device__ __forceinline__ T minimum(T a, T b) {
    return (isnan(b) || b < a) ? b : a;
}

template <typename scalar_t>
__global__ void min_softmax_kernel(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int depth,
    int64_t spatial
) {
    using acc_t = at::acc_type<scalar_t, true>;

    __shared__ acc_t partial_max[8][32];
    __shared__ acc_t partial_sum[8][32];

    const int lane = threadIdx.x;
    const int group = threadIdx.y;
    const int64_t position = int64_t(blockIdx.x) * 32 + lane;
    const int64_t batch = blockIdx.y;
    const bool valid = position < spatial;

    acc_t m0 = acc_t(INFINITY);
    acc_t m1 = acc_t(INFINITY);
    acc_t m2 = acc_t(INFINITY);

    if (valid) {
        const int64_t channel_stride = int64_t(depth) * spatial;
        const scalar_t* p =
            input + (batch * 24 + group) * channel_stride + position;

        #pragma unroll 1
        for (int d = 0; d < depth; ++d) {
            const int64_t offset = int64_t(d) * spatial;
            m0 = minimum(m0, acc_t(p[offset]));
            m1 = minimum(m1, acc_t(p[8 * channel_stride + offset]));
            m2 = minimum(m2, acc_t(p[16 * channel_stride + offset]));
        }
    }

    acc_t local_max = m0 > m1 ? m0 : m1;
    local_max = local_max > m2 ? local_max : m2;
    partial_max[group][lane] = local_max;
    __syncthreads();

    acc_t maximum = partial_max[0][lane];
    #pragma unroll
    for (int i = 1; i < 8; ++i) {
        const acc_t value = partial_max[i][lane];
        maximum = maximum > value ? maximum : value;
    }

    const acc_t e0 = exponential(m0 - maximum);
    const acc_t e1 = exponential(m1 - maximum);
    const acc_t e2 = exponential(m2 - maximum);

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

    acc_t denominator = partial_sum[0][lane];
    #pragma unroll
    for (int i = 1; i < 8; ++i) {
        denominator += partial_sum[i][lane];
    }

    if (valid) {
        const int64_t base = (batch * 24 + group) * spatial + position;
        output[base] = scalar_t(e0 / denominator);
        output[base + 8 * spatial] = scalar_t(e1 / denominator);
        output[base + 16 * spatial] = scalar_t(e2 / denominator);
    }
}

torch::Tensor min_softmax_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(x.dim() == 5 && x.size(1) == 24,
                "Expected an NCDHW tensor with 24 channels");
    TORCH_CHECK(x.is_contiguous(), "Expected a contiguous tensor");
    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 spatial = x.size(3) * x.size(4);

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

    const dim3 threads(32, 8);
    const dim3 blocks(
        static_cast<unsigned int>((spatial + 31) / 32),
        static_cast<unsigned int>(batch)
    );
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        x.scalar_type(),
        "min_softmax_cuda",
        [&] {
            min_softmax_kernel<scalar_t><<<blocks, threads, 0, stream>>>(
                x.data_ptr<scalar_t>(),
                output.data_ptr<scalar_t>(),
                depth,
                spatial
            );
        }
    );
    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=["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 self.dim in (2, -3)
            and x.size(1) == 24
            and x.size(0) <= 65535
            and x.is_contiguous()
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
        ):
            return _fused_ops.min_softmax_cuda(x)
        return torch.softmax(torch.min(x, dim=self.dim)[0], dim=1)
```