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

_CPP = """
#include <torch/extension.h>

torch::Tensor conv1d_tf32_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias);
"""

_CUDA = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;

// Each block computes 32 output channels by 128 output positions.
// Splitting operands into two TF32 values substantially improves accuracy
// over a single TF32 multiplication.
__global__ void conv1d_tf32_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int output_channels,
    int length,
    int output_length)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    constexpr int BM = 32;
    constexpr int BN = 128;
    constexpr int BK = 32;

    __shared__ __align__(32) float ah[BM * BK];
    __shared__ __align__(32) float al[BM * BK];
    __shared__ __align__(32) float bh[BK * BN];
    __shared__ __align__(32) float bl[BK * BN];

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int warp_row = (warp / 4) * 16;
    const int warp_col = (warp % 4) * 32;

    const int position_base = blockIdx.x * BN;
    const int channel_base = blockIdx.y * BM;
    const int batch = blockIdx.z;
    const int reduction_size = channels * 3;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[2];
    wmma::fill_fragment(acc[0], 0.0f);
    wmma::fill_fragment(acc[1], 0.0f);

    for (int base = 0; base < reduction_size; base += BK) {
        for (int i = tid; i < BM * BK; i += blockDim.x) {
            const int oc = channel_base + i / BK;
            const int k = base + i % BK;
            float value = 0.0f;
            if (oc < output_channels && k < reduction_size) {
                value = weight[(int64_t)oc * reduction_size + k];
            }
            const float hi = wmma::__float_to_tf32(value);
            ah[i] = hi;
            al[i] = wmma::__float_to_tf32(value - hi);
        }

        for (int i = tid; i < BK * BN; i += blockDim.x) {
            const int k = base + i / BN;
            const int position = position_base + i % BN;
            float value = 0.0f;
            if (k < reduction_size && position < output_length) {
                const int ic = k / 3;
                const int offset = k % 3;
                value = x[
                    ((int64_t)batch * channels + ic) * length
                    + position + offset];
            }
            const float hi = wmma::__float_to_tf32(value);
            bh[i] = hi;
            bl[i] = wmma::__float_to_tf32(value - hi);
        }
        __syncthreads();

        #pragma unroll
        for (int k = 0; k < BK; k += 8) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> a_hi, a_lo;

            wmma::load_matrix_sync(a_hi, ah + warp_row * BK + k, BK);
            wmma::load_matrix_sync(a_lo, al + warp_row * BK + k, BK);

            #pragma unroll
            for (int tile = 0; tile < 2; ++tile) {
                wmma::fragment<
                    wmma::matrix_b, 16, 16, 8,
                    wmma::precision::tf32, wmma::row_major> b_hi, b_lo;

                const int offset = k * BN + warp_col + tile * 16;
                wmma::load_matrix_sync(b_hi, bh + offset, BN);
                wmma::load_matrix_sync(b_lo, bl + offset, BN);

                wmma::mma_sync(acc[tile], a_lo, b_hi, acc[tile]);
                wmma::mma_sync(acc[tile], a_hi, b_lo, acc[tile]);
                wmma::mma_sync(acc[tile], a_hi, b_hi, acc[tile]);
            }
        }
        __syncthreads();
    }

    // Reuse the input tile storage for the output tile.
    #pragma unroll
    for (int tile = 0; tile < 2; ++tile) {
        wmma::store_matrix_sync(
            bh + warp_row * BN + warp_col + tile * 16,
            acc[tile], BN, wmma::mem_row_major);
    }
    __syncthreads();

    for (int i = tid; i < BM * BN; i += blockDim.x) {
        const int oc = channel_base + i / BN;
        const int position = position_base + i % BN;
        if (oc < output_channels && position < output_length) {
            float value = bh[i];
            if (bias != nullptr) {
                value += bias[oc];
            }
            output[
                ((int64_t)batch * output_channels + oc) * output_length
                + position] = value;
        }
    }
#endif
}

torch::Tensor conv1d_tf32_cuda(
    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(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Input and weight must be float32");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Input and weight must be contiguous");
    TORCH_CHECK(x.dim() == 3 && weight.dim() == 3 &&
                weight.size(2) == 3 && x.size(1) == weight.size(1),
                "Invalid convolution dimensions");
    TORCH_CHECK(x.device() == weight.device(),
                "Input and weight must be on the same device");
    TORCH_CHECK(x.size(2) >= 3, "Input length must be at least three");

    const int batch = x.size(0);
    const int channels = x.size(1);
    const int output_channels = weight.size(0);
    const int length = x.size(2);
    const int output_length = length - 2;

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

    auto output = torch::empty(
        {batch, output_channels, output_length}, x.options());

    if (output.numel() == 0) {
        return output;
    }

    const dim3 grid(
        (output_length + 127) / 128,
        (output_channels + 31) / 32,
        batch);

    conv1d_tf32_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        channels,
        output_channels,
        length,
        output_length);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="conv1d_split_tf32_k3",
            cpp_sources=_CPP,
            cuda_sources=_CUDA,
            functions=["conv1d_tf32_cuda"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _extension


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

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv1d
        supported = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 3
            and x.is_contiguous()
            and conv.weight.is_cuda
            and conv.weight.dtype == torch.float32
            and conv.weight.is_contiguous()
            and conv.kernel_size == (3,)
            and conv.stride == (1,)
            and conv.padding == (0,)
            and conv.dilation == (1,)
            and conv.groups == 1
            and x.shape[2] >= 3
            and x.shape[0] <= 65535
            and not torch.is_autocast_enabled()
        )
        if supported and torch.cuda.get_device_capability(x.device)[0] >= 8:
            return _get_extension().conv1d_tf32_cuda(
                x, conv.weight, conv.bias
            )
        return conv(x)
```