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


_cuda_source = r"""
#include <torch/extension.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;

constexpr int NTILE = 128;
constexpr int KTILE = 32;

__global__ __launch_bounds__(128)
void asymmetric_conv_tf32(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int height,
    int width,
    int depth,
    int output_width,
    int spatial_size)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    __shared__ __align__(32) float storage[12288];

    float* ah = storage;
    float* al = ah + 64 * KTILE;
    float* bh = al + 64 * KTILE;
    float* bl = bh + KTILE * NTILE;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int batch = blockIdx.y;
    const int start = blockIdx.x * NTILE;
    const int position = start + tid;

    const int row_stride = output_width * depth;
    const int ih = position / row_stride;
    const int rem = position - ih * row_stride;
    const int iw = rem / depth;
    const int iz = rem - iw * depth;

    const int64_t channel_stride =
        static_cast<int64_t>(height) * width * depth;
    const int64_t base =
        static_cast<int64_t>(batch) * 3 * channel_stride
        + (static_cast<int64_t>(ih) * width + iw) * depth + iz;

    #pragma unroll
    for (int ic = 0; ic < 3; ++ic) {
        #pragma unroll
        for (int kh = 0; kh < 3; ++kh) {
            #pragma unroll
            for (int kw = 0; kw < 3; ++kw) {
                const int k = ic * 9 + kh * 3 + kw;
                float v = 0.0f;
                if (position < spatial_size) {
                    v = input[
                        base + ic * channel_stride
                        + static_cast<int64_t>(kh) * width * depth
                        + kw * depth];
                }
                const float hi = wmma::__float_to_tf32(v);
                bh[k * NTILE + tid] = hi;
                bl[k * NTILE + tid] = wmma::__float_to_tf32(v - hi);
            }
        }
    }

    #pragma unroll
    for (int k = 27; k < KTILE; ++k) {
        bh[k * NTILE + tid] = 0.0f;
        bl[k * NTILE + tid] = 0.0f;
    }

    #pragma unroll
    for (int i = tid; i < 64 * KTILE; i += 128) {
        const int oc = i / KTILE;
        const int k = i % KTILE;
        const float v = k < 27 ? weight[oc * 27 + k] : 0.0f;
        const float hi = wmma::__float_to_tf32(v);
        ah[i] = hi;
        al[i] = wmma::__float_to_tf32(v - hi);
    }

    __syncthreads();

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[8];

    #pragma unroll
    for (int j = 0; j < 8; ++j) {
        wmma::fill_fragment(acc[j], 0.0f);
    }

    #pragma unroll 1
    for (int k = 0; k < KTILE; k += 8) {
        wmma::fragment<
            wmma::matrix_a, 16, 16, 8,
            wmma::precision::tf32, wmma::row_major> a_hi, a_lo;

        wmma::load_matrix_sync(a_hi, ah + warp * 16 * KTILE + k, KTILE);
        wmma::load_matrix_sync(a_lo, al + warp * 16 * KTILE + k, KTILE);

        #pragma unroll
        for (int j = 0; j < 8; ++j) {
            wmma::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> b_hi, b_lo;

            wmma::load_matrix_sync(b_hi, bh + k * NTILE + j * 16, NTILE);
            wmma::load_matrix_sync(b_lo, bl + k * NTILE + j * 16, NTILE);

            // Three products recover FP32 accuracy from TF32 components.
            wmma::mma_sync(acc[j], a_hi, b_hi, acc[j]);
            wmma::mma_sync(acc[j], a_lo, b_hi, acc[j]);
            wmma::mma_sync(acc[j], a_hi, b_lo, acc[j]);
        }
    }

    // Reuse the operand storage for coalesced output writes.
    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 8; ++j) {
        wmma::store_matrix_sync(
            storage + warp * 16 * NTILE + j * 16,
            acc[j], NTILE, wmma::mem_row_major);
    }

    __syncthreads();

    if ((spatial_size & 3) == 0) {
        #pragma unroll
        for (int i = tid; i < 64 * NTILE / 4; i += 128) {
            const int oc = i / (NTILE / 4);
            const int col = (i % (NTILE / 4)) * 4;
            const int pos = start + col;

            if (pos < spatial_size) {
                float4 v = reinterpret_cast<const float4*>(storage)[i];
                if (bias != nullptr) {
                    const float b = bias[oc];
                    v.x += b;
                    v.y += b;
                    v.z += b;
                    v.w += b;
                }

                float* dst = output
                    + (static_cast<int64_t>(batch) * 64 + oc) * spatial_size
                    + pos;
                *reinterpret_cast<float4*>(dst) = v;
            }
        }
    } else {
        #pragma unroll
        for (int i = tid; i < 64 * NTILE; i += 128) {
            const int oc = i / NTILE;
            const int pos = start + i % NTILE;
            if (pos < spatial_size) {
                float v = storage[i];
                if (bias != nullptr) v += bias[oc];
                output[
                    (static_cast<int64_t>(batch) * 64 + oc) * spatial_size
                    + pos] = v;
            }
        }
    }
#endif
}

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias)
{
    const c10::cuda::CUDAGuard guard(input.device());

    const int batch = input.size(0);
    const int height = input.size(2);
    const int width = input.size(3);
    const int depth = input.size(4);
    const int output_height = height - 2;
    const int output_width = width - 2;
    const int spatial_size = output_height * output_width * depth;

    auto output = torch::empty(
        {batch, 64, output_height, output_width, depth},
        input.options());

    if (output.numel() == 0) return output;

    const dim3 grid((spatial_size + NTILE - 1) / NTILE, batch);
    const cudaStream_t stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();

    asymmetric_conv_tf32<<<grid, 128, 0, stream>>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        height, width, depth, output_width, spatial_size);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_cpp_source = r"""
#include <torch/extension.h>
torch::Tensor asymmetric_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias);
"""

_extension = load_inline(
    name="asymmetric_conv3d_tf32_components_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["asymmetric_conv_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: int,
        stride: int = 1,
        padding: int = 0,
        dilation: int = 1,
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv3d = nn.Conv3d(
            in_channels,
            out_channels,
            (kernel_size, kernel_size, 1),
            stride=stride,
            padding=padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv3d
        weight = conv.weight
        bias = conv.bias

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.is_contiguous()
            and x.shape[1] == 3
            and x.shape[2] >= 3
            and x.shape[3] >= 3
            and x.shape[4] > 0
            and x.shape[0] <= 65535
            and x.shape[2] * x.shape[3] * x.shape[4] < 2147483392
            and weight.device == x.device
            and weight.dtype == torch.float32
            and weight.is_contiguous()
            and conv.in_channels == 3
            and conv.out_channels == 64
            and conv.kernel_size == (3, 3, 1)
            and conv.stride == (1, 1, 1)
            and conv.padding == (0, 0, 0)
            and conv.dilation == (1, 1, 1)
            and conv.groups == 1
            and not torch.is_autocast_enabled()
            and torch.cuda.get_device_capability(x.device)[0] >= 8
            and (
                bias is None
                or (
                    bias.device == x.device
                    and bias.dtype == torch.float32
                    and bias.is_contiguous()
                )
            )
        ):
            return _extension.asymmetric_conv_cuda(
                x, weight, weight if bias is None else bias, bias is not None
            )

        return conv(x)
```