```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 fast_conv1d(
    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>
#include <mma.h>

namespace {
using namespace nvcuda;

constexpr int BM = 128;
constexpr int BN = 128;
constexpr int BK = 32;
constexpr int AS = 40;
constexpr int BS = 136;
constexpr int SHMEM_BYTES = BM * BN * sizeof(float);

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

    half* ah = reinterpret_cast<half*>(storage);
    half* al = ah + BM * AS;
    half* bh = al + BM * AS;
    half* bl = bh + BK * BS;

    const int tid = threadIdx.x;
    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 + (int64_t)blockIdx.y * 64 * length;

    using AFrag = wmma::fragment<
        wmma::matrix_a, 16, 16, 16, half, wmma::row_major>;
    using BFrag = wmma::fragment<
        wmma::matrix_b, 16, 16, 16, half, wmma::row_major>;
    using CFrag = wmma::fragment<
        wmma::accumulator, 16, 16, 16, float>;

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

    #pragma unroll 1
    for (int base = 0; base < 192; base += BK) {
        #pragma unroll
        for (int index = tid; index < BM * BK; index += 256) {
            const int row = index / BK;
            const int k = index % BK;
            const float value = w[row * 192 + base + k];
            const half high = __float2half_rn(value);
            ah[row * AS + k] = high;
            al[row * AS + k] =
                __float2half_rn(value - __half2float(high));
        }

        #pragma unroll
        for (int index = tid; index < BK * BN; index += 256) {
            const int k = index / BN;
            const int col = index % BN;
            const int reduction = base + k;
            const int channel = reduction / 3;
            const int tap = reduction - channel * 3;
            const int position = start + col;

            float value = 0.0f;
            if (position < output_length)
                value = input[(int64_t)channel * length + position + tap];

            const half high = __float2half_rn(value);
            bh[k * BS + col] = high;
            bl[k * BS + col] =
                __float2half_rn(value - __half2float(high));
        }

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < BK; k += 16) {
            AFrag a_high[2], a_low[2];

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                wmma::load_matrix_sync(
                    a_high[i], ah + (wm + i * 16) * AS + k, AS);
                wmma::load_matrix_sync(
                    a_low[i], al + (wm + i * 16) * AS + k, AS);
            }

            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                BFrag b_high, b_low;
                wmma::load_matrix_sync(
                    b_high, bh + k * BS + wn + j * 16, BS);
                wmma::load_matrix_sync(
                    b_low, bl + k * BS + wn + j * 16, BS);

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

        __syncthreads();
    }

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

    #pragma unroll
    for (int i = 0; i < 2; ++i) {
        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            wmma::store_matrix_sync(
                tile + (wm + i * 16) * BN + wn + j * 16,
                acc[i][j], BN, wmma::mem_row_major);
        }
    }

    __syncthreads();

    if ((output_length & 1) == 0) {
        #pragma unroll
        for (int index = tid; index < BM * BN / 2; index += 256) {
            const int row = index / (BN / 2);
            const int col = (index % (BN / 2)) * 2;
            if (start + col + 1 < output_length) {
                float2 value = reinterpret_cast<float2*>(tile)[index];
                const float b = bias ? bias[row] : 0.0f;
                value.x += b;
                value.y += b;

                const int64_t offset =
                    ((int64_t)blockIdx.y * 128 + row) * output_length
                    + start + col;
                *reinterpret_cast<float2*>(out + offset) = value;
            }
        }
    } else {
        #pragma unroll
        for (int index = tid; index < BM * BN; index += 256) {
            const int row = index / BN;
            const int col = index % BN;
            if (start + col < output_length) {
                const int64_t offset =
                    ((int64_t)blockIdx.y * 128 + row) * output_length
                    + start + col;
                out[offset] = tile[index] + (bias ? bias[row] : 0.0f);
            }
        }
    }
#endif
}
}

torch::Tensor fast_conv1d(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    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,
                "Invalid input shape");
    TORCH_CHECK(weight.dim() == 3 && weight.size(0) == 128 &&
                weight.size(1) == 64 && weight.size(2) == 3,
                "Invalid weight shape");

    const int64_t batch = x.size(0);
    const int64_t length = x.size(2);
    TORCH_CHECK(batch <= 65535 && length <= 2147483647,
                "Input exceeds kernel indexing limits");

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

    if (batch == 0)
        return output;

    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.numel() == 128,
                    "Invalid bias");
        bias_ptr = b.data_ptr<float>();
    }

    C10_CUDA_CHECK(cudaFuncSetAttribute(
        conv_kernel,
        cudaFuncAttributeMaxDynamicSharedMemorySize,
        SHMEM_BYTES));

    const dim3 grid((output_length + BN - 1) / BN, batch);
    conv_kernel<<<grid, 256, SHMEM_BYTES,
                  at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        static_cast<int>(length),
        output_length);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="conv1d_tensorcore_split_fp32_v2",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fast_conv1d"],
    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.fast_conv1d(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 x.dim() == 3
            and x.size(1) == 64
            and 3 <= x.size(2) <= 2147483647
            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:
            properties = torch.cuda.get_device_properties(device)
            self._supported_devices[device] = (
                properties.major >= 7
                and properties.shared_memory_per_block_optin >= 65536
            )

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

        return _FastConv1d.apply(
            x.contiguous(), conv.weight.contiguous(), conv.bias
        )
```