```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_runtime.h>
#include <math.h>

template<int FIXED_DEPTH>
__global__ void depth_min_softmax_kernel(
    const float* __restrict__ x,
    float* __restrict__ out,
    int depth,
    int spatial)
{
    const int lane = threadIdx.x;
    const int group = threadIdx.y;
    const int pixel = blockIdx.x * 32 + lane;
    const int batch = blockIdx.y;
    const bool valid = pixel < spatial;

    const int D = FIXED_DEPTH > 0 ? FIXED_DEPTH : depth;
    const int64_t channel_stride = static_cast<int64_t>(D) * spatial;
    const int64_t base =
        static_cast<int64_t>(batch) * 24 * channel_stride + pixel;

    float minima[3] = {INFINITY, INFINITY, INFINITY};

    if (valid) {
        #pragma unroll
        for (int d = 0; d < D; ++d) {
            #pragma unroll
            for (int k = 0; k < 3; ++k) {
                const int channel = group + 8 * k;
                const float value =
                    x[base + channel * channel_stride +
                      static_cast<int64_t>(d) * spatial];
                minima[k] =
                    (value < minima[k] || isnan(value))
                    ? value : minima[k];
            }
        }
    }

    __shared__ float scratch[8][32];

    scratch[group][lane] =
        fmaxf(minima[0], fmaxf(minima[1], minima[2]));
    __syncthreads();

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

    float exponentials[3];
    float local_sum = 0.0f;
    #pragma unroll
    for (int k = 0; k < 3; ++k) {
        exponentials[k] = __expf(minima[k] - maximum);
        local_sum += exponentials[k];
    }

    // All warps must finish reading maxima before reusing scratch.
    __syncthreads();
    scratch[group][lane] = local_sum;
    __syncthreads();

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

    if (valid) {
        const float inverse_sum = 1.0f / denominator;
        #pragma unroll
        for (int k = 0; k < 3; ++k) {
            const int channel = group + 8 * k;
            const int64_t index =
                (static_cast<int64_t>(batch) * 24 + channel) *
                spatial + pixel;
            out[index] = exponentials[k] * inverse_sum;
        }
    }
}

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 && x.size(1) == 24,
                "Expected shape [N, 24, D, H, W]");
    TORCH_CHECK(x.is_contiguous(), "Expected contiguous input");
    TORCH_CHECK(x.size(2) > 0, "Depth must be positive");

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

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

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

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

    const dim3 threads(32, 8);
    const dim3 blocks((spatial + 31) / 32, batch);
    const auto stream = at::cuda::getCurrentCUDAStream();

    if (depth == 22) {
        depth_min_softmax_kernel<22><<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), out.data_ptr<float>(), depth, spatial);
    } else {
        depth_min_softmax_kernel<0><<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), out.data_ptr<float>(), depth, spatial);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_fused_ops = load_inline(
    name="conv3d_depth_min_softmax_fused_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.is_contiguous()
            and x.shape[1] == 24
            and self.dim in (2, -3)
        ):
            return _fused_ops.depth_min_softmax_cuda(x)
        return torch.softmax(torch.min(x, dim=self.dim)[0], dim=1)
```