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

constexpr int TILE = 128;

union __align__(16) SharedStorage {
    struct {
        unsigned ah[16][72];
        unsigned al[16][72];
        unsigned bh[16][136];
        unsigned bl[16][136];
    } matrix;
    float output[64][136];
};

__device__ __forceinline__
unsigned pack_halves(__half a, __half b) {
    return __half2_as_uint(__halves2half2(a, b));
}

__device__ __forceinline__
void split_pair(float x, float y, unsigned& hi, unsigned& lo) {
    const __half hx = __float2half_rn(x);
    const __half hy = __float2half_rn(y);
    hi = pack_halves(hx, hy);
    lo = pack_halves(
        __float2half_rn(x - __half2float(hx)),
        __float2half_rn(y - __half2float(hy)));
}

__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
}

__global__ __launch_bounds__(128)
void asymmetric_conv_kernel(
    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__ SharedStorage shared;

    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;

    const int batch = blockIdx.y;
    const int start = blockIdx.x * TILE;
    const int position = start + tid;
    const int input_row_stride = width * depth;
    const int output_row_stride = output_width * depth;
    const int ih = position / output_row_stride;
    const int remainder = position - ih * output_row_stride;

    const int64_t channel_stride =
        static_cast<int64_t>(height) * input_row_stride;
    const int64_t base =
        static_cast<int64_t>(batch) * 3 * channel_stride
        + static_cast<int64_t>(ih) * input_row_stride + remainder;

    #pragma unroll
    for (int p = 0; p < 16; ++p) {
        constexpr int kernel_elements = 27;
        const int k0 = 2 * p;
        const int k1 = k0 + 1;
        float x0 = 0.0f;
        float x1 = 0.0f;

        if (position < spatial_size) {
            if (k0 < kernel_elements) {
                x0 = input[
                    base + (k0 / 9) * channel_stride
                    + ((k0 % 9) / 3) * input_row_stride
                    + (k0 % 3) * depth];
            }
            if (k1 < kernel_elements) {
                x1 = input[
                    base + (k1 / 9) * channel_stride
                    + ((k1 % 9) / 3) * input_row_stride
                    + (k1 % 3) * depth];
            }
        }

        unsigned hi, lo;
        split_pair(x0, x1, hi, lo);
        shared.matrix.bh[p][tid] = hi;
        shared.matrix.bl[p][tid] = lo;
    }

    #pragma unroll
    for (int i = tid; i < 16 * 64; i += 128) {
        const int p = i >> 6;
        const int oc = i & 63;
        const int k = p * 2;
        const float w0 = k < 27 ? weight[oc * 27 + k] : 0.0f;
        const float w1 = k + 1 < 27 ? weight[oc * 27 + k + 1] : 0.0f;

        unsigned hi, lo;
        split_pair(w0, w1, hi, lo);
        shared.matrix.ah[p][oc] = hi;
        shared.matrix.al[p][oc] = lo;
    }

    __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 < 16; k += 8) {
        const int p = k + pair;

        const unsigned a0 = shared.matrix.ah[p][row];
        const unsigned a1 = shared.matrix.ah[p][row + 8];
        const unsigned a2 = shared.matrix.ah[p + 4][row];
        const unsigned a3 = shared.matrix.ah[p + 4][row + 8];

        const unsigned l0 = shared.matrix.al[p][row];
        const unsigned l1 = shared.matrix.al[p][row + 8];
        const unsigned l2 = shared.matrix.al[p + 4][row];
        const unsigned l3 = shared.matrix.al[p + 4][row + 8];

        #pragma unroll
        for (int j = 0; j < 16; ++j) {
            const int col = j * 8 + group;
            const unsigned b0 = shared.matrix.bh[p][col];
            const unsigned b1 = shared.matrix.bh[p + 4][col];
            const unsigned r0 = shared.matrix.bl[p][col];
            const unsigned r1 = shared.matrix.bl[p + 4][col];

            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);
        }
    }

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

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 16; ++j) {
        const int col = j * 8 + pair * 2;
        *reinterpret_cast<float2*>(&shared.output[row][col]) =
            make_float2(acc[j][0] + bias0, acc[j][1] + bias0);
        *reinterpret_cast<float2*>(&shared.output[row + 8][col]) =
            make_float2(acc[j][2] + bias1, acc[j][3] + bias1);
    }

    __syncthreads();

    const int pos = start + lane * 4;

    #pragma unroll
    for (int c = 0; c < 16; ++c) {
        const int oc = warp + c * 4;
        const float4 value =
            *reinterpret_cast<const float4*>(&shared.output[oc][lane * 4]);

        float* dst = output
            + (static_cast<int64_t>(batch) * 64 + oc) * spatial_size;

        if ((spatial_size & 3) == 0) {
            if (pos < spatial_size) {
                *reinterpret_cast<float4*>(dst + pos) = value;
            }
        } else {
            if (pos < spatial_size) dst[pos] = value.x;
            if (pos + 1 < spatial_size) dst[pos + 1] = value.y;
            if (pos + 2 < spatial_size) dst[pos + 2] = value.z;
            if (pos + 3 < spatial_size) dst[pos + 3] = value.w;
        }
    }
#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_kernel<<<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_packed_mma_v5",
    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)
```