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


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

torch::Tensor to_channels_last_cuda(torch::Tensor x);
torch::Tensor depth_min_bias_softmax_cuda(
    torch::Tensor x, torch::Tensor bias);
"""

_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_constants.h>

__global__ void interleave_three_channels(
    const float* __restrict__ input,
    float4* __restrict__ output,
    int spatial,
    int vectors)
{
    const int v = blockIdx.x * blockDim.x + threadIdx.x;
    if (v >= vectors) return;

    const int64_t batch_offset = int64_t(blockIdx.y) * spatial * 3;
    const float* src = input + batch_offset;
    const int q = v * 4;

    float4 result;
    result.x = src[int64_t(q % 3) * spatial + q / 3];
    result.y = src[int64_t((q + 1) % 3) * spatial + (q + 1) / 3];
    result.z = src[int64_t((q + 2) % 3) * spatial + (q + 2) / 3];
    result.w = src[int64_t((q + 3) % 3) * spatial + (q + 3) / 3];

    output[int64_t(blockIdx.y) * vectors + v] = result;
}

torch::Tensor to_channels_last_cuda(torch::Tensor x)
{
    TORCH_CHECK(
        x.is_cuda() && x.scalar_type() == at::kFloat &&
        x.dim() == 5 && x.size(1) == 3 && x.is_contiguous(),
        "Expected contiguous float32 CUDA input with three channels");

    const int64_t spatial64 = x.size(2) * x.size(3) * x.size(4);
    TORCH_CHECK(
        spatial64 % 4 == 0 && spatial64 <= 715827880LL &&
        x.size(0) <= 65535,
        "Unsupported input dimensions");

    const c10::cuda::CUDAGuard guard(x.device());
    auto storage = torch::empty(
        {x.size(0), x.size(2), x.size(3), x.size(4), 3},
        x.options());

    if (storage.numel() != 0) {
        const int spatial = static_cast<int>(spatial64);
        const int vectors = spatial * 3 / 4;
        const dim3 grid(
            (int64_t(vectors) + 255) / 256,
            static_cast<unsigned int>(x.size(0)));

        interleave_three_channels<<<
            grid, 256, 0, at::cuda::getCurrentCUDAStream()
        >>>(
            x.data_ptr<float>(),
            reinterpret_cast<float4*>(storage.data_ptr<float>()),
            spatial,
            vectors);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return storage.permute({0, 4, 1, 2, 3});
}

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

__device__ __forceinline__ float4 minimum4(float4 a, float4 b)
{
    return make_float4(
        minimum_nan(a.x, b.x),
        minimum_nan(a.y, b.y),
        minimum_nan(a.z, b.z),
        minimum_nan(a.w, b.w));
}

template<int STATIC_DEPTH>
__global__ void depth_min_bias_softmax_kernel(
    const float4* __restrict__ input,
    const float4* __restrict__ bias,
    float4* __restrict__ output,
    int spatial,
    int dynamic_depth)
{
    constexpr int GROUP_SIZE = 8;
    constexpr int PIXELS_PER_BLOCK = 32;
    constexpr int CHANNEL_VECTORS = 6;

    const int channel_vector = threadIdx.x & 7;
    const int position =
        blockIdx.x * PIXELS_PER_BLOCK + (threadIdx.x >> 3);
    const int depth = STATIC_DEPTH ? STATIC_DEPTH : dynamic_depth;
    const bool active =
        position < spatial && channel_vector < CHANNEL_VECTORS;

    float4 minimum = make_float4(
        CUDART_INF_F, CUDART_INF_F, CUDART_INF_F, CUDART_INF_F);
    float4 second = minimum;

    if (active) {
        const int64_t stride = int64_t(spatial) * CHANNEL_VECTORS;
        const int64_t base =
            int64_t(blockIdx.y) * depth * stride +
            int64_t(position) * CHANNEL_VECTORS + channel_vector;

        if (STATIC_DEPTH == 22) {
            #pragma unroll
            for (int d = 0; d < 11; ++d) {
                const float4 a = input[base + int64_t(d) * stride];
                const float4 b = input[base + int64_t(d + 11) * stride];
                minimum = minimum4(minimum, a);
                second = minimum4(second, b);
            }
            minimum = minimum4(minimum, second);
        } else {
            for (int d = 0; d < depth; ++d) {
                minimum = minimum4(
                    minimum, input[base + int64_t(d) * stride]);
            }
        }

        const float4 b = bias[channel_vector];
        minimum.x += b.x;
        minimum.y += b.y;
        minimum.z += b.z;
        minimum.w += b.w;
    }

    float maximum = active
        ? fmaxf(fmaxf(minimum.x, minimum.y), fmaxf(minimum.z, minimum.w))
        : -CUDART_INF_F;

    #pragma unroll
    for (int offset = 4; offset > 0; offset >>= 1) {
        maximum = fmaxf(
            maximum,
            __shfl_xor_sync(0xffffffff, maximum, offset, GROUP_SIZE));
    }

    float4 e = make_float4(0.f, 0.f, 0.f, 0.f);
    if (active) {
        e.x = __expf(minimum.x - maximum);
        e.y = __expf(minimum.y - maximum);
        e.z = __expf(minimum.z - maximum);
        e.w = __expf(minimum.w - maximum);
    }

    float sum = (e.x + e.y) + (e.z + e.w);
    #pragma unroll
    for (int offset = 4; offset > 0; offset >>= 1) {
        sum += __shfl_xor_sync(0xffffffff, sum, offset, GROUP_SIZE);
    }

    if (active) {
        const float inverse = 1.f / sum;
        e.x *= inverse;
        e.y *= inverse;
        e.z *= inverse;
        e.w *= inverse;

        output[
            (int64_t(blockIdx.y) * spatial + position) *
            CHANNEL_VECTORS + channel_vector
        ] = e;
    }
}

torch::Tensor depth_min_bias_softmax_cuda(
    torch::Tensor x, torch::Tensor bias)
{
    TORCH_CHECK(
        x.is_cuda() && bias.is_cuda() &&
        x.scalar_type() == at::kFloat &&
        bias.scalar_type() == at::kFloat,
        "Expected float32 CUDA tensors");
    TORCH_CHECK(
        x.dim() == 5 && x.size(1) == 24 && x.size(2) > 0 &&
        x.is_contiguous(at::MemoryFormat::ChannelsLast3d),
        "Expected channels-last tensor of shape (N, 24, D, H, W)");
    TORCH_CHECK(
        bias.numel() == 24 && bias.is_contiguous() &&
        bias.device() == x.device(),
        "Invalid bias");
    TORCH_CHECK(x.size(0) <= 65535, "Batch exceeds launch limit");

    const c10::cuda::CUDAGuard guard(x.device());
    const int64_t spatial64 = x.size(3) * x.size(4);
    TORCH_CHECK(
        spatial64 <= 2147483616LL && x.size(2) <= 2147483647LL,
        "Tensor dimensions are too large");

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

    const int spatial = static_cast<int>(spatial64);
    const int depth = static_cast<int>(x.size(2));
    const dim3 grid(
        static_cast<unsigned int>((spatial64 + 31) / 32),
        static_cast<unsigned int>(x.size(0)));
    const auto stream = at::cuda::getCurrentCUDAStream();

    const auto* input_ptr =
        reinterpret_cast<const float4*>(x.data_ptr<float>());
    const auto* bias_ptr =
        reinterpret_cast<const float4*>(bias.data_ptr<float>());
    auto* output_ptr =
        reinterpret_cast<float4*>(storage.data_ptr<float>());

    if (depth == 22) {
        depth_min_bias_softmax_kernel<22><<<grid, 256, 0, stream>>>(
            input_ptr, bias_ptr, output_ptr, spatial, depth);
    } else {
        depth_min_bias_softmax_kernel<0><<<grid, 256, 0, stream>>>(
            input_ptr, bias_ptr, output_ptr, spatial, depth);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="conv3d_min_softmax_interleave_v6",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=[
        "to_channels_last_cuda",
        "depth_min_bias_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.conv.to(memory_format=torch.channels_last_3d)
        self.dim = dim

    def forward(self, x):
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.shape[0] <= 65535
            and self.conv.weight.dtype == torch.float32
            and self.conv.out_channels == 24
            and self.conv.bias is not None
            and self.conv.bias.dtype == torch.float32
            and self.dim in (2, -3)
        ):
            if not x.is_contiguous(memory_format=torch.channels_last_3d):
                spatial = x.shape[2] * x.shape[3] * x.shape[4]
                if (
                    x.is_contiguous()
                    and x.shape[1] == 3
                    and spatial % 4 == 0
                    and spatial <= 715827880
                ):
                    x = _fused_ops.to_channels_last_cuda(x)
                else:
                    x = x.contiguous(memory_format=torch.channels_last_3d)

            y = F.conv3d(
                x,
                self.conv.weight,
                None,
                self.conv.stride,
                self.conv.padding,
                self.conv.dilation,
                self.conv.groups,
            )

            if y.dtype == torch.float32:
                return _fused_ops.depth_min_bias_softmax_cuda(
                    y.contiguous(memory_format=torch.channels_last_3d),
                    self.conv.bias,
                )

            y = y + self.conv.bias.to(y.dtype).view(1, -1, 1, 1, 1)
            return torch.softmax(torch.min(y, dim=self.dim)[0], dim=1)

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