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

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

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

    __shared__ acc_t partial[8][32];

    const int lane = threadIdx.x;
    const int group = threadIdx.y;
    const int64_t batch = blockIdx.x / tiles;
    const int64_t position =
        (static_cast<int64_t>(blockIdx.x) % tiles) * 32 + lane;

    acc_t minima[3] = {
        static_cast<acc_t>(INFINITY),
        static_cast<acc_t>(INFINITY),
        static_cast<acc_t>(INFINITY)
    };

    if (position < spatial) {
        const int64_t base =
            (batch * 24 + group) * depth * spatial + position;

        for (int64_t d = 0; d < depth; ++d) {
            #pragma unroll
            for (int i = 0; i < 3; ++i) {
                const acc_t value = static_cast<acc_t>(
                    input[base + (static_cast<int64_t>(i) * 8 * depth + d)
                                     * spatial]);
                if (isnan(value) || value < minima[i]) {
                    minima[i] = value;
                }
            }
        }
    }

    acc_t local_max = static_cast<acc_t>(-INFINITY);
    #pragma unroll
    for (int i = 0; i < 3; ++i) {
        if (minima[i] > local_max) {
            local_max = minima[i];
        }
    }
    partial[group][lane] = local_max;
    __syncthreads();

    acc_t maximum = static_cast<acc_t>(-INFINITY);
    #pragma unroll
    for (int g = 0; g < 8; ++g) {
        const acc_t value = partial[g][lane];
        if (value > maximum) {
            maximum = value;
        }
    }
    // Finish every shared-memory read before reusing partial.
    __syncthreads();

    acc_t numerators[3];
    acc_t local_sum = static_cast<acc_t>(0);
    #pragma unroll
    for (int i = 0; i < 3; ++i) {
        numerators[i] = exp(minima[i] - maximum);
        local_sum += numerators[i];
    }
    partial[group][lane] = local_sum;
    __syncthreads();

    acc_t denominator = static_cast<acc_t>(0);
    #pragma unroll
    for (int g = 0; g < 8; ++g) {
        denominator += partial[g][lane];
    }

    if (position < spatial) {
        #pragma unroll
        for (int i = 0; i < 3; ++i) {
            const int channel = group + i * 8;
            output[(batch * 24 + channel) * spatial + position] =
                static_cast<scalar_t>(numerators[i] / denominator);
        }
    }
}

torch::Tensor min_depth_softmax_cuda(torch::Tensor input)
{
    TORCH_CHECK(input.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(input.dim() == 5, "Expected NCDHW input");
    TORCH_CHECK(input.is_contiguous(), "Expected contiguous input");
    TORCH_CHECK(input.size(1) == 24, "Expected 24 channels");
    TORCH_CHECK(input.size(2) > 0, "Depth must be positive");

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

    const int64_t batch = input.size(0);
    const int64_t depth = input.size(2);
    const int64_t height = input.size(3);
    const int64_t width = input.size(4);
    const int64_t spatial = height * width;

    auto output = torch::empty({batch, 24, height, width}, input.options());
    if (output.numel() == 0) {
        return output;
    }

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

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        input.scalar_type(),
        "min_depth_softmax_cuda",
        [&] {
            min_depth_softmax_kernel<scalar_t>
                <<<blocks, threads, 0, stream>>>(
                    input.data_ptr<scalar_t>(),
                    output.data_ptr<scalar_t>(),
                    depth, spatial, tiles);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="fused_conv3d_depth_min_softmax_v1",
    cpp_sources=(
        "torch::Tensor min_depth_softmax_cuda(torch::Tensor input);"
    ),
    cuda_sources=_cuda_source,
    functions=["min_depth_softmax_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _MinDepthSoftmax(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        output = _extension.min_depth_softmax_cuda(x.contiguous())
        if ctx.needs_input_grad[0]:
            ctx.save_for_backward(x, output)
        return output

    @staticmethod
    def backward(ctx, grad_output):
        x, output = ctx.saved_tensors
        indices = x.min(dim=2).indices
        grad_min = output * (
            grad_output - (grad_output * output).sum(dim=1, keepdim=True)
        )
        grad_x = torch.zeros_like(x)
        grad_x.scatter_(2, indices.unsqueeze(2), grad_min.unsqueeze(2))
        return grad_x


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.shape[1] == 24
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
        ):
            return _MinDepthSoftmax.apply(x)
        return torch.softmax(torch.min(x, dim=self.dim)[0], dim=1)
```