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

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

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

template<int STATIC_DEPTH>
__global__ __launch_bounds__(256)
void depth_min_bias_softmax_kernel(
    const float4* __restrict__ input,
    const float4* __restrict__ bias,
    float4* __restrict__ output,
    int spatial,
    int dynamic_depth)
{
    constexpr int THREADS = 256;
    constexpr int TILE_PIXELS = 128;
    constexpr int CHANNEL_VECTORS = 6;
    constexpr int VECTORS_PER_THREAD = 3;

    __shared__ float4 reduced[TILE_PIXELS * CHANNEL_VECTORS];

    const int tid = threadIdx.x;
    const int depth = STATIC_DEPTH ? STATIC_DEPTH : dynamic_depth;
    const int64_t depth_stride = int64_t(spatial) * CHANNEL_VECTORS;
    const int64_t tile_start =
        int64_t(blockIdx.x) * TILE_PIXELS * CHANNEL_VECTORS;
    const int64_t batch_start =
        int64_t(blockIdx.y) * depth * depth_stride;

    float4 minima[VECTORS_PER_THREAD];
    bool valid[VECTORS_PER_THREAD];

    #pragma unroll
    for (int k = 0; k < VECTORS_PER_THREAD; ++k) {
        minima[k] = make_float4(
            CUDART_INF_F, CUDART_INF_F, CUDART_INF_F, CUDART_INF_F);
        valid[k] = tile_start + tid + k * THREADS < depth_stride;
    }

    // Three independent, fully coalesced streams per thread.
    #pragma unroll 2
    for (int d = 0; d < depth; ++d) {
        const int64_t base =
            batch_start + int64_t(d) * depth_stride + tile_start + tid;
        #pragma unroll
        for (int k = 0; k < VECTORS_PER_THREAD; ++k) {
            if (valid[k]) {
                minima[k] = minimum4(
                    minima[k], input[base + k * THREADS]);
            }
        }
    }

    #pragma unroll
    for (int k = 0; k < VECTORS_PER_THREAD; ++k) {
        reduced[tid + k * THREADS] = minima[k];
    }
    __syncthreads();

    const int channel_vector = tid & 7;
    const float4 b = channel_vector < CHANNEL_VECTORS
        ? bias[channel_vector]
        : make_float4(0.f, 0.f, 0.f, 0.f);

    #pragma unroll
    for (int iteration = 0; iteration < 4; ++iteration) {
        const int local_pixel = (tid >> 3) + iteration * 32;
        const int64_t pixel =
            int64_t(blockIdx.x) * TILE_PIXELS + local_pixel;
        const bool active =
            pixel < spatial && channel_vector < CHANNEL_VECTORS;

        float4 value = make_float4(0.f, 0.f, 0.f, 0.f);
        float maximum = -CUDART_INF_F;

        if (active) {
            value = reduced[local_pixel * CHANNEL_VECTORS + channel_vector];
            value.x += b.x;
            value.y += b.y;
            value.z += b.z;
            value.w += b.w;
            maximum = fmaxf(
                fmaxf(value.x, value.y), fmaxf(value.z, value.w));
        }

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

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

        float denominator =
            (exponential.x + exponential.y) +
            (exponential.z + exponential.w);

        #pragma unroll
        for (int offset = 4; offset > 0; offset >>= 1) {
            denominator +=
                __shfl_xor_sync(0xffffffff, denominator, offset, 8);
        }

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

            const int64_t index =
                (int64_t(blockIdx.y) * spatial + pixel) *
                CHANNEL_VECTORS + channel_vector;
            output[index] = exponential;
        }
    }
}

torch::Tensor depth_min_bias_softmax_cuda(
    torch::Tensor x, torch::Tensor bias)
{
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(
        x.scalar_type() == at::kFloat &&
        bias.scalar_type() == at::kFloat,
        "Expected float32 tensors");
    TORCH_CHECK(
        x.dim() == 5 && x.size(1) == 24 && x.size(2) > 0,
        "Expected shape (N, 24, D, H, W)");
    TORCH_CHECK(
        x.is_contiguous(at::MemoryFormat::ChannelsLast3d),
        "Expected channels-last 3D storage");
    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 <= 2147483647LL && 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 + 127) / 128),
        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_bias_softmax_coalesced_v5",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["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 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)
            and x.ndim == 5
            and x.shape[0] <= 65535
        ):
            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
        )
```