```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.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <stdint.h>

constexpr int TILE = 128;
constexpr int KS = 40;

__device__ __forceinline__
void mma_f16(
    float* d,
    unsigned a0, unsigned a1, unsigned a2, unsigned a3,
    unsigned b0, unsigned b1)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    asm volatile(
        "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 "
        "{%0, %1, %2, %3}, "
        "{%4, %5, %6, %7}, "
        "{%8, %9}, "
        "{%0, %1, %2, %3};\n"
        : "+f"(d[0]), "+f"(d[1]), "+f"(d[2]), "+f"(d[3])
        : "r"(a0), "r"(a1), "r"(a2), "r"(a3),
          "r"(b0), "r"(b1));
#endif
}

__device__ __forceinline__
unsigned load_pair(const __half* p)
{
    return *reinterpret_cast<const unsigned*>(p);
}

__device__ __forceinline__
void write_pair(float* dst, int pos, int size, float x, float y)
{
    if ((size & 1) == 0) {
        if (pos < size) {
            *reinterpret_cast<float2*>(dst + pos) = make_float2(x, y);
        }
    } else {
        if (pos < size) dst[pos] = x;
        if (pos + 1 < size) dst[pos + 1] = y;
    }
}

__global__ __launch_bounds__(128)
void asymmetric_conv_mma(
    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__(16) __half ah[64][KS];
    __shared__ __align__(16) __half al[64][KS];
    __shared__ __align__(16) __half bh[TILE][KS];
    __shared__ __align__(16) __half bl[TILE][KS];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int warp = tid >> 5;
    const int group = lane >> 2;
    const int pair = (lane & 3) * 2;

    const int batch = blockIdx.y;
    const int start = blockIdx.x * TILE;
    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 __half hi = __float2half_rn(v);
                bh[tid][k] = hi;
                bl[tid][k] = __float2half_rn(v - __half2float(hi));
            }
        }
    }

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

    #pragma unroll
    for (int i = tid; i < 64 * 32; i += 128) {
        const int oc = i >> 5;
        const int k = i & 31;
        const float v = k < 27 ? weight[oc * 27 + k] : 0.0f;
        const __half hi = __float2half_rn(v);
        ah[oc][k] = hi;
        al[oc][k] = __float2half_rn(v - __half2float(hi));
    }

    __syncthreads();

    float acc[16][4];

    #pragma unroll
    for (int j = 0; j < 16; ++j) {
        #pragma unroll
        for (int q = 0; q < 4; ++q) {
            acc[j][q] = 0.0f;
        }
    }

    const int row = warp * 16 + group;

    #pragma unroll
    for (int k = 0; k < 32; k += 16) {
        const unsigned a0 = load_pair(&ah[row][k + pair]);
        const unsigned a1 = load_pair(&ah[row + 8][k + pair]);
        const unsigned a2 = load_pair(&ah[row][k + pair + 8]);
        const unsigned a3 = load_pair(&ah[row + 8][k + pair + 8]);

        const unsigned l0 = load_pair(&al[row][k + pair]);
        const unsigned l1 = load_pair(&al[row + 8][k + pair]);
        const unsigned l2 = load_pair(&al[row][k + pair + 8]);
        const unsigned l3 = load_pair(&al[row + 8][k + pair + 8]);

        #pragma unroll
        for (int j = 0; j < 16; ++j) {
            const int col = j * 8 + group;
            const unsigned b0 = load_pair(&bh[col][k + pair]);
            const unsigned b1 = load_pair(&bh[col][k + pair + 8]);
            const unsigned r0 = load_pair(&bl[col][k + pair]);
            const unsigned r1 = load_pair(&bl[col][k + pair + 8]);

            mma_f16(acc[j], a0, a1, a2, a3, b0, b1);
            mma_f16(acc[j], l0, l1, l2, l3, b0, b1);
            mma_f16(acc[j], a0, a1, a2, a3, r0, r1);
        }
    }

    float* dst0 = output
        + (static_cast<int64_t>(batch) * 64 + row) * spatial_size;
    float* dst1 = dst0 + static_cast<int64_t>(8) * spatial_size;

    const float bias0 = bias ? bias[row] : 0.0f;
    const float bias1 = bias ? bias[row + 8] : 0.0f;

    #pragma unroll
    for (int j = 0; j < 16; ++j) {
        const int pos = start + j * 8 + pair;
        write_pair(
            dst0, pos, spatial_size,
            acc[j][0] + bias0, acc[j][1] + bias0);
        write_pair(
            dst1, pos, spatial_size,
            acc[j][2] + bias1, acc[j][3] + bias1);
    }
#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 + TILE - 1) / TILE, batch);
    const cudaStream_t stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();

    asymmetric_conv_mma<<<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_mma_components_v4",
    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)
```