```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 BM = 128;
constexpr int BN = 128;
constexpr int AS = 24;
constexpr int BS = 18;
constexpr int OS = 136;
constexpr int A_ELEMENTS = 3 * BM * AS;
constexpr int B_ELEMENTS = (BN + 2) * BS;
constexpr int SHARED_BYTES =
    2 * (A_ELEMENTS + B_ELEMENTS) * sizeof(half);

__device__ __forceinline__
unsigned shared_address(const void* ptr) {
    return static_cast<unsigned>(__cvta_generic_to_shared(ptr));
}

__device__ __forceinline__
void load_a(unsigned (&r)[4], const half* ptr) {
#if __CUDA_ARCH__ >= 800
    const unsigned address = shared_address(ptr);
    asm volatile(
        "ldmatrix.sync.aligned.m8n8.x4.shared.b16 "
        "{%0, %1, %2, %3}, [%4];"
        : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3])
        : "r"(address));
#endif
}

__device__ __forceinline__
void mma(float (&c)[4], const unsigned (&a)[4],
         const unsigned (&b)[2]) {
#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[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]),
          "r"(b[0]), "r"(b[1]));
#endif
}

__global__ void pack_weights(
    const float* __restrict__ weight,
    half* __restrict__ packed)
{
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= 4 * 3 * 128 * 8)
        return;

    const int pair = index & 7;
    const int row = (index >> 3) & 127;
    const int group = index >> 10;
    const int tap = group % 3;
    const int stage = group / 3;
    const int channel = stage * 16 + pair * 2;

    const float v0 = weight[row * 192 + channel * 3 + tap];
    const float v1 = weight[row * 192 + (channel + 1) * 3 + tap];

    const half2 high = __floats2half2_rn(v0, v1);
    const float2 rounded = __half22float2(high);
    const half2 low = __floats2half2_rn(
        v0 - rounded.x, v1 - rounded.y);

    const int offset = group * 4096 + row * 16 + pair * 2;
    *reinterpret_cast<half2*>(packed + offset) = high;
    *reinterpret_cast<half2*>(packed + offset + 2048) = low;
}

__device__ __forceinline__
void store_input(
    half* high,
    half* low,
    int channel,
    int position,
    float value)
{
    const half h = __float2half_rn(value);
    high[position * BS + channel] = h;
    low[position * BS + channel] =
        __float2half_rn(value - __half2float(h));
}

__global__ __launch_bounds__(256, 2)
void conv_kernel(
    const float* __restrict__ x,
    const half* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int length,
    int output_length)
{
#if __CUDA_ARCH__ >= 800
    extern __shared__ __align__(32) unsigned char storage[];

    half* ah = reinterpret_cast<half*>(storage);
    half* al = ah + A_ELEMENTS;
    half* bh = al + A_ELEMENTS;
    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) {
        // Each input value is converted once and reused by all three taps.
        #pragma unroll
        for (int p = tid; p < 16 * BN; p += 256) {
            const int channel = p / BN;
            const int position = p % BN;
            const int global_position = start + position;

            float value = 0.0f;
            if (global_position < length) {
                value = input[
                    static_cast<int64_t>(stage * 16 + channel) * length
                    + global_position];
            }

            store_input(bh, bl, channel, position, value);
        }

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

            float value = 0.0f;
            if (global_position < length) {
                value = input[
                    static_cast<int64_t>(stage * 16 + channel) * length
                    + global_position];
            }

            store_input(bh, bl, channel, position, value);
        }

        #pragma unroll
        for (int tap = 0; tap < 3; ++tap) {
            const int row = tid >> 1;
            const int channel = (tid & 1) * 8;
            const int src =
                (stage * 3 + tap) * 4096 + row * 16 + channel;
            const int dst = tap * BM * AS + row * AS + channel;

            *reinterpret_cast<uint4*>(ah + dst) =
                *reinterpret_cast<const uint4*>(weight + src);
            *reinterpret_cast<uint4*>(al + dst) =
                *reinterpret_cast<const uint4*>(weight + src + 2048);
        }

        __syncthreads();

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

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                const int offset =
                    tap * BM * AS
                    + (wm + i * 16 + (lane & 15)) * AS
                    + (lane >> 4) * 8;

                load_a(a_high[i], ah + offset);
                load_a(a_low[i], al + offset);
            }

            #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;

                unsigned b_high[2];
                unsigned b_low[2];

                b_high[0] =
                    *reinterpret_cast<const unsigned*>(bh + offset);
                b_high[1] =
                    *reinterpret_cast<const unsigned*>(bh + offset + 8);
                b_low[0] =
                    *reinterpret_cast<const unsigned*>(bl + offset);
                b_low[1] =
                    *reinterpret_cast<const unsigned*>(bl + offset + 8);

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

        __syncthreads();
    }

    float* tile = reinterpret_cast<float*>(storage);

    #pragma unroll
    for (int section = 0; section < 2; ++section) {
        if ((warp >> 2) == section) {
            const int local_m = wm - section * 64;
            const int fragment_row = lane >> 2;
            const int fragment_col = (lane & 3) * 2;

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                const int row = local_m + i * 16 + fragment_row;

                #pragma unroll
                for (int j = 0; j < 8; ++j) {
                    const int col = wn + j * 8 + fragment_col;

                    *reinterpret_cast<float2*>(tile + row * OS + col) =
                        make_float2(accum[i][j][0], accum[i][j][1]);

                    *reinterpret_cast<float2*>(
                        tile + (row + 8) * OS + col) =
                        make_float2(accum[i][j][2], accum[i][j][3]);
                }
            }
        }

        __syncthreads();

        #pragma unroll
        for (int p = tid; p < 64 * BN / 2; p += 256) {
            const int row = p / (BN / 2);
            const int col = (p % (BN / 2)) * 2;
            const int channel = section * 64 + row;
            const int position = start + col;

            if (position < output_length) {
                float2 value = *reinterpret_cast<const float2*>(
                    tile + row * OS + col);

                if (bias) {
                    const float b = bias[channel];
                    value.x += b;
                    value.y += b;
                }

                const int64_t offset =
                    (static_cast<int64_t>(blockIdx.y) * 128 + channel)
                    * output_length + position;

                if ((output_length & 1) == 0) {
                    *reinterpret_cast<float2*>(output + offset) = value;
                } else {
                    output[offset] = value.x;
                    if (position + 1 < output_length)
                        output[offset + 1] = value.y;
                }
            }
        }

        if (section == 0)
            __syncthreads();
    }
#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");

    const int batch = static_cast<int>(x.size(0));
    const int length = static_cast<int>(x.size(2));
    const int output_length = length - 2;
    auto output = torch::empty({batch, 128, output_length}, x.options());

    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(
        {4 * 3 * 2 * 128 * 16},
        weight.options().dtype(at::kHalf));

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

    pack_weights<<<48, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        reinterpret_cast<half*>(packed.data_ptr<at::Half>()));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    const dim3 grid((output_length + BN - 1) / BN, batch);
    conv_kernel<<<grid, 256, SHARED_BYTES, stream>>>(
        x.data_ptr<float>(),
        reinterpret_cast<const half*>(packed.data_ptr<at::Half>()),
        bias_ptr,
        output.data_ptr<float>(),
        length,
        output_length);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_extension = load_inline(
    name="conv1d_reuse_input_split_fp32_v4",
    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 not torch.is_autocast_enabled()
        )

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