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

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

torch::Tensor conv1d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<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.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>

namespace {

constexpr int BN = 128;
constexpr int BS = 18;
constexpr int THREADS = 256;
constexpr int B_ELEMENTS = (BN + 2) * BS;
constexpr int SHARED_BYTES = 2 * B_ELEMENTS * sizeof(half);

__device__ __forceinline__
void split_pair(float a, float b, unsigned& high, unsigned& low) {
    const half2 h = __floats2half2_rn(a, b);
    const float2 f = __half22float2(h);
    high = __half2_as_uint(h);
    low = __half2_as_uint(
        __floats2half2_rn(a - f.x, b - f.y));
}

__global__ void pack_weights(
    const float* __restrict__ weight,
    uint4* __restrict__ packed) {
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= 12 * 8 * 32)
        return;

    const int lane = index & 31;
    const int tile = index >> 5;
    const int group = tile >> 3;
    const int row = (tile & 7) * 16 + (lane >> 2);
    const int channel = (group / 3) * 16 + (lane & 3) * 2;
    const int tap = group % 3;

    const int p = row * 192 + channel * 3 + tap;

    uint4 high, low;
    split_pair(
        weight[p], weight[p + 3],
        high.x, low.x);
    split_pair(
        weight[p + 8 * 192], weight[p + 8 * 192 + 3],
        high.y, low.y);
    split_pair(
        weight[p + 8 * 3], weight[p + 8 * 3 + 3],
        high.z, low.z);
    split_pair(
        weight[p + 8 * 192 + 8 * 3],
        weight[p + 8 * 192 + 8 * 3 + 3],
        high.w, low.w);

    packed[tile * 64 + lane] = high;
    packed[tile * 64 + 32 + lane] = low;
}

__device__ __forceinline__
float load_input(const float* ptr) {
    float result;
    asm volatile(
        "ld.global.cg.f32 %0, [%1];"
        : "=f"(result)
        : "l"(ptr));
    return result;
}

__device__ __forceinline__
float2 load_input_pair(const float* ptr) {
    float2 result;
    asm volatile(
        "ld.global.cg.v2.f32 {%0, %1}, [%2];"
        : "=f"(result.x), "=f"(result.y)
        : "l"(ptr));
    return result;
}

__device__ __forceinline__
uint4 load_weight(const uint4* ptr) {
    uint4 result;
    asm volatile(
        "ld.global.ca.v4.u32 {%0, %1, %2, %3}, [%4];"
        : "=r"(result.x), "=r"(result.y),
          "=r"(result.z), "=r"(result.w)
        : "l"(ptr));
    return result;
}

__device__ __forceinline__
void store_input_pair(
    half* high,
    half* low,
    int offset,
    float a,
    float b) {
    const half2 h = __floats2half2_rn(a, b);
    const float2 rounded = __half22float2(h);

    *reinterpret_cast<half2*>(high + offset) = h;
    *reinterpret_cast<half2*>(low + offset) =
        __floats2half2_rn(a - rounded.x, b - rounded.y);
}

__device__ __forceinline__
void mma(
    float (&c)[4],
    const uint4& a,
    unsigned b0,
    unsigned b1) {
#if __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};"
        : "+f"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
        : "r"(a.x), "r"(a.y), "r"(a.z), "r"(a.w),
          "r"(b0), "r"(b1));
#endif
}

__global__ __launch_bounds__(THREADS, 2)
void conv_kernel(
    const float* __restrict__ x,
    const uint4* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int length,
    int output_length,
    int output_stride) {
#if __CUDA_ARCH__ >= 800
    extern __shared__ __align__(32) unsigned char storage[];
    half* bh = reinterpret_cast<half*>(storage);
    half* bl = bh + B_ELEMENTS;

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int warp = tid >> 5;
    const int wm = (warp >> 1) * 32;
    const int wn = (warp & 1) * 64;
    const int start = blockIdx.x * BN;

    const float* input =
        x + static_cast<int64_t>(blockIdx.y) * 64 * length;

    float accum[2][8][4];

    #pragma unroll
    for (int i = 0; i < 2; ++i)
        #pragma unroll
        for (int j = 0; j < 8; ++j)
            #pragma unroll
            for (int r = 0; r < 4; ++r)
                accum[i][j][r] = 0.0f;

    #pragma unroll 1
    for (int stage = 0; stage < 4; ++stage) {
        #pragma unroll
        for (int p = tid; p < 8 * (BN / 2); p += THREADS) {
            const int channel = (p / (BN / 2)) * 2;
            const int position = (p % (BN / 2)) * 2;
            const int global_position = start + position;

            const float* p0 =
                input
                + static_cast<int64_t>(stage * 16 + channel) * length;
            const float* p1 = p0 + length;

            float2 v0 = make_float2(0.0f, 0.0f);
            float2 v1 = make_float2(0.0f, 0.0f);

            if ((length & 1) == 0 && global_position + 1 < length) {
                v0 = load_input_pair(p0 + global_position);
                v1 = load_input_pair(p1 + global_position);
            } else {
                if (global_position < length) {
                    v0.x = load_input(p0 + global_position);
                    v1.x = load_input(p1 + global_position);
                }
                if (global_position + 1 < length) {
                    v0.y = load_input(p0 + global_position + 1);
                    v1.y = load_input(p1 + global_position + 1);
                }
            }

            store_input_pair(
                bh, bl, position * BS + channel, v0.x, v1.x);
            store_input_pair(
                bh, bl, (position + 1) * BS + channel, v0.y, v1.y);
        }

        if (tid < 16) {
            const int channel = (tid >> 1) * 2;
            const int position = BN + (tid & 1);
            const int global_position = start + position;

            float v0 = 0.0f;
            float v1 = 0.0f;
            if (global_position < length) {
                const float* ptr =
                    input
                    + static_cast<int64_t>(stage * 16 + channel) * length
                    + global_position;
                v0 = load_input(ptr);
                v1 = load_input(ptr + length);
            }
            store_input_pair(
                bh, bl, position * BS + channel, v0, v1);
        }

        __syncthreads();

        #pragma unroll
        for (int tap = 0; tap < 3; ++tap) {
            uint4 a_high[2];
            uint4 a_low[2];

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                const uint4* ptr =
                    weight
                    + ((stage * 3 + tap) * 8 + wm / 16 + i) * 64
                    + lane;

                a_high[i] = load_weight(ptr);
                a_low[i] = load_weight(ptr + 32);
            }

            #pragma unroll
            for (int j = 0; j < 8; ++j) {
                const int position =
                    wn + j * 8 + (lane >> 2) + tap;
                const int channel = (lane & 3) * 2;
                const int offset = position * BS + channel;

                const unsigned b_high0 =
                    *reinterpret_cast<const unsigned*>(bh + offset);
                const unsigned b_high1 =
                    *reinterpret_cast<const unsigned*>(bh + offset + 8);
                const unsigned b_low0 =
                    *reinterpret_cast<const unsigned*>(bl + offset);
                const unsigned b_low1 =
                    *reinterpret_cast<const unsigned*>(bl + offset + 8);

                #pragma unroll
                for (int i = 0; i < 2; ++i) {
                    mma(accum[i][j], a_high[i], b_low0, b_low1);
                    mma(accum[i][j], a_low[i], b_high0, b_high1);
                    mma(accum[i][j], a_high[i], b_high0, b_high1);
                }
            }
        }

        if (stage != 3)
            __syncthreads();
    }

    // Padded channel strides allow fully utilized 32-byte store sectors.
    #pragma unroll
    for (int i = 0; i < 2; ++i) {
        const int row = wm + i * 16 + (lane >> 2);
        const float bias0 = bias ? bias[row] : 0.0f;
        const float bias1 = bias ? bias[row + 8] : 0.0f;

        float* out0 =
            output
            + (static_cast<int64_t>(blockIdx.y) * 128 + row)
              * output_stride;
        float* out1 = out0 + static_cast<int64_t>(8) * output_stride;

        #pragma unroll
        for (int j = 0; j < 8; ++j) {
            const int position =
                start + wn + j * 8 + (lane & 3) * 2;

            if (position + 1 < output_length) {
                *reinterpret_cast<float2*>(out0 + position) =
                    make_float2(
                        accum[i][j][0] + bias0,
                        accum[i][j][1] + bias0);
                *reinterpret_cast<float2*>(out1 + position) =
                    make_float2(
                        accum[i][j][2] + bias1,
                        accum[i][j][3] + bias1);
            } else if (position < output_length) {
                out0[position] = accum[i][j][0] + bias0;
                out1[position] = accum[i][j][2] + bias1;
            }
        }
    }
#endif
}

} // namespace

torch::Tensor conv1d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(
        x.scalar_type() == at::kFloat &&
        weight.scalar_type() == at::kFloat,
        "Expected float32 tensors");
    TORCH_CHECK(
        x.is_contiguous() && weight.is_contiguous(),
        "Expected contiguous tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(
        x.dim() == 3 && x.size(1) == 64 &&
        x.size(2) >= 3 && x.size(2) <= 2147483520LL &&
        x.size(0) <= 65535,
        "Unsupported input shape");
    TORCH_CHECK(
        weight.dim() == 3 &&
        weight.size(0) == 128 &&
        weight.size(1) == 64 &&
        weight.size(2) == 3,
        "Unsupported weight shape");
    TORCH_CHECK(
        at::cuda::getCurrentDeviceProperties()->major >= 8,
        "Requires compute capability 8.0 or newer");

    const int batch = static_cast<int>(x.size(0));
    const int length = static_cast<int>(x.size(2));
    const int output_length = length - 2;
    const int output_stride = (output_length + 31) & ~31;

    auto storage = torch::empty(
        {batch, 128, output_stride}, x.options());
    auto output = storage.narrow(2, 0, output_length);

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(
            b.device() == x.device() &&
            b.scalar_type() == at::kFloat &&
            b.is_contiguous() &&
            b.dim() == 1 && b.numel() == 128,
            "Invalid bias");
        bias_ptr = b.data_ptr<float>();
    }

    if (batch == 0)
        return output;

    auto packed = torch::empty(
        {12 * 8 * 64 * 4}, weight.options().dtype(at::kInt));

    const auto stream = at::cuda::getCurrentCUDAStream();

    pack_weights<<<12, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        reinterpret_cast<uint4*>(packed.data_ptr<int>()));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    const dim3 grid((output_length + BN - 1) / BN, batch);
    conv_kernel<<<grid, THREADS, SHARED_BYTES, stream>>>(
        x.data_ptr<float>(),
        reinterpret_cast<const uint4*>(packed.data_ptr<int>()),
        bias_ptr,
        output.data_ptr<float>(),
        length,
        output_length,
        output_stride);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

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


class _FastConv1d(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias):
        ctx.save_for_backward(x, weight)
        ctx.has_bias = bias is not None
        return _extension.conv1d_cuda(x, weight, bias)

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        grad_output = grad_output.contiguous()
        grad_x = grad_weight = grad_bias = None

        if ctx.needs_input_grad[0]:
            grad_x = torch.nn.grad.conv1d_input(
                x.shape, weight, grad_output
            )
        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv1d_weight(
                x, weight.shape, grad_output
            )
        if ctx.has_bias and ctx.needs_input_grad[2]:
            grad_bias = grad_output.sum(dim=(0, 2))

        return grad_x, grad_weight, grad_bias


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.conv1d = nn.Conv1d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )
        self._supported_devices = {}

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv1d

        eligible = (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and x.dim() == 3
            and x.size(1) == 64
            and 3 <= x.size(2) <= 2147483520
            and x.size(0) <= 65535
            and conv.in_channels == 64
            and conv.out_channels == 128
            and conv.kernel_size == (3,)
            and conv.stride == (1,)
            and conv.padding == (0,)
            and conv.dilation == (1,)
            and conv.groups == 1
            and conv.padding_mode == "zeros"
            and not torch.is_autocast_enabled()
            and (
                conv.bias is None
                or (
                    conv.bias.dtype == torch.float32
                    and conv.bias.device == x.device
                )
            )
        )

        if not eligible:
            return conv(x)

        device = x.device.index
        if device not in self._supported_devices:
            self._supported_devices[device] = (
                torch.cuda.get_device_capability(device)[0] >= 8
            )

        if not self._supported_devices[device]:
            return conv(x)

        return _FastConv1d.apply(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias.contiguous() if conv.bias is not None else None,
        )
```