```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_tf32_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 <mma.h>

namespace {

constexpr int TILE_M = 64;
constexpr int TILE_N = 64;
constexpr int TILE_K = 32;
constexpr int REDUCTION = 64 * 3;

union __align__(32) SharedStorage {
    struct {
        float ah[TILE_M * TILE_K];
        float al[TILE_M * TILE_K];
        float bh[TILE_K * TILE_N];
        float bl[TILE_K * TILE_N];
    } inputs;
    float output[TILE_M * TILE_N];
};

__global__ void conv1d_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int length,
    int output_length)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    __shared__ SharedStorage shared;

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int warp_row = (warp / 2) * 32;
    const int warp_col = (warp % 2) * 32;
    const int first_channel = blockIdx.y * TILE_M;
    const int first_position = blockIdx.x * TILE_N;
    const int batch = blockIdx.z;

    const float* batch_x =
        x + static_cast<int64_t>(batch) * 64 * length;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[2][2];

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

    for (int base_k = 0; base_k < REDUCTION; base_k += TILE_K) {
        #pragma unroll
        for (int index = tid; index < TILE_M * TILE_K; index += 128) {
            const int row = index / TILE_K;
            const int k = base_k + index % TILE_K;
            const float value =
                weight[(first_channel + row) * REDUCTION + k];
            const float high = __float_to_tf32(value);

            shared.inputs.ah[index] = high;
            shared.inputs.al[index] = __float_to_tf32(value - high);
        }

        #pragma unroll
        for (int index = tid; index < TILE_K * TILE_N; index += 128) {
            const int k = base_k + index / TILE_N;
            const int position = first_position + index % TILE_N;
            const int input_channel = k / 3;
            const int kernel_position = k % 3;

            float value = 0.0f;
            if (position < output_length) {
                value = batch_x[
                    static_cast<int64_t>(input_channel) * length +
                    position + kernel_position];
            }

            const float high = __float_to_tf32(value);
            shared.inputs.bh[index] = high;
            shared.inputs.bl[index] = __float_to_tf32(value - high);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < TILE_K; kk += 8) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> ah[2], al[2];
            wmma::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> bh[2], bl[2];

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                const int offset = (warp_row + 16 * i) * TILE_K + kk;
                wmma::load_matrix_sync(
                    ah[i], shared.inputs.ah + offset, TILE_K);
                wmma::load_matrix_sync(
                    al[i], shared.inputs.al + offset, TILE_K);
            }

            #pragma unroll
            for (int j = 0; j < 2; ++j) {
                const int offset = kk * TILE_N + warp_col + 16 * j;
                wmma::load_matrix_sync(
                    bh[j], shared.inputs.bh + offset, TILE_N);
                wmma::load_matrix_sync(
                    bl[j], shared.inputs.bl + offset, TILE_N);
            }

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                #pragma unroll
                for (int j = 0; j < 2; ++j) {
                    // Three products retain substantially more precision
                    // than a single TF32 multiplication.
                    wmma::mma_sync(acc[i][j], al[i], bh[j], acc[i][j]);
                    wmma::mma_sync(acc[i][j], ah[i], bl[j], acc[i][j]);
                    wmma::mma_sync(acc[i][j], ah[i], bh[j], acc[i][j]);
                }
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 2; ++i) {
        #pragma unroll
        for (int j = 0; j < 2; ++j) {
            wmma::store_matrix_sync(
                shared.output +
                    (warp_row + 16 * i) * TILE_N + warp_col + 16 * j,
                acc[i][j],
                TILE_N,
                wmma::mem_row_major);
        }
    }

    __syncthreads();

    for (int index = tid; index < TILE_M * TILE_N; index += 128) {
        const int channel = first_channel + index / TILE_N;
        const int position = first_position + index % TILE_N;

        if (position < output_length) {
            float value = shared.output[index];
            if (bias != nullptr) {
                value += bias[channel];
            }
            output[
                (static_cast<int64_t>(batch) * 128 + channel) *
                    output_length + position] = value;
        }
    }
#endif
}

} // namespace

torch::Tensor conv1d_tf32_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    const c10::cuda::CUDAGuard device_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.dim() == 3 && x.size(1) == 64 && x.size(2) >= 3,
        "Expected input shape [N, 64, L], L >= 3");
    TORCH_CHECK(
        weight.dim() == 3 && weight.size(0) == 128 &&
        weight.size(1) == 64 && weight.size(2) == 3,
        "Expected weight shape [128, 64, 3]");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(
        at::cuda::getCurrentDeviceProperties()->major >= 8,
        "This kernel requires compute capability 8.0 or newer");

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

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

    if (x.size(0) == 0) {
        return output;
    }

    const dim3 grid(
        (output_length + TILE_N - 1) / TILE_N,
        128 / TILE_M,
        static_cast<unsigned int>(x.size(0)));

    conv1d_kernel<<<grid, 128, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        length,
        output_length);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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


class _Conv1dFunction(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_tf32_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

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 3
            and x.shape[1] == 64
            and 3 <= x.shape[2] < 2**31 - 64
            and x.shape[0] <= 65535
            and x.is_contiguous()
            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.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            and not torch.is_autocast_enabled()
        ):
            device = x.device.index
            if device not in self._supported_devices:
                self._supported_devices[device] = (
                    torch.cuda.get_device_capability(device)[0] >= 8
                )

            if self._supported_devices[device]:
                return _Conv1dFunction.apply(x, conv.weight, conv.bias)

        return conv(x)
```