```python
import torch
import torch.nn as nn
import torch.nn.functional as F
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,
    int64_t stride,
    int64_t padding,
    int64_t dilation,
    int64_t groups);
"""

_cuda_source = 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>

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
__device__ __forceinline__ float round_tf32(float x) {
    float result;
    asm("cvt.rn.tf32.f32 %0, %1;" : "=f"(result) : "f"(x));
    return result;
}
#endif

__global__ void conv1d_tf32_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int64_t length,
    int64_t output_length,
    int input_channels,
    int output_channels,
    int channels_per_group,
    int outputs_per_group,
    int kernel_size,
    int stride,
    int padding,
    int dilation,
    int groups) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    // Shared storage is reused for the output after the reduction.
    __shared__ __align__(32) float storage[4096];
    float* shared_a = storage;
    float* shared_b = storage + 2048;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_row = (warp >> 1) * 32;
    const int warp_col = (warp & 1) * 32;

    const int group = blockIdx.z % groups;
    const int batch = blockIdx.z / groups;
    const int output_channel_start = blockIdx.y * 64;
    const int64_t position_start = int64_t(blockIdx.x) * 64;
    const int reduction_size = channels_per_group * kernel_size;

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

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

    for (int reduction_start = 0;
         reduction_start < reduction_size;
         reduction_start += 32) {

        for (int index = tid; index < 2048; index += 128) {
            const int row = index / 32;
            const int k = reduction_start + index % 32;
            const int local_output_channel = output_channel_start + row;

            float value = 0.0f;
            if (local_output_channel < outputs_per_group &&
                k < reduction_size) {
                const int output_channel =
                    group * outputs_per_group + local_output_channel;
                value = weight[
                    int64_t(output_channel) * reduction_size + k];
            }
            shared_a[index] = value;
        }

        for (int index = tid; index < 2048; index += 128) {
            const int k = reduction_start + index / 64;
            const int64_t position = position_start + index % 64;

            float value = 0.0f;
            if (k < reduction_size && position < output_length) {
                const int channel = k / kernel_size;
                const int tap = k % kernel_size;
                const int64_t input_position =
                    position * stride - padding + int64_t(tap) * dilation;

                if (input_position >= 0 && input_position < length) {
                    const int input_channel =
                        group * channels_per_group + channel;
                    value = x[
                        (int64_t(batch) * input_channels + input_channel)
                        * length + input_position];
                }
            }
            shared_b[index] = value;
        }

        __syncthreads();

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

            wmma::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> b_hi[2], b_lo[2];

            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                wmma::load_matrix_sync(
                    a_hi[i], shared_a + (warp_row + i * 16) * 32 + k, 32);

                #pragma unroll
                for (int element = 0;
                     element < a_hi[i].num_elements; ++element) {
                    const float original = a_hi[i].x[element];
                    const float high = round_tf32(original);
                    a_hi[i].x[element] = high;
                    a_lo[i].x[element] = round_tf32(original - high);
                }
            }

            #pragma unroll
            for (int j = 0; j < 2; ++j) {
                wmma::load_matrix_sync(
                    b_hi[j], shared_b + k * 64 + warp_col + j * 16, 64);

                #pragma unroll
                for (int element = 0;
                     element < b_hi[j].num_elements; ++element) {
                    const float original = b_hi[j].x[element];
                    const float high = round_tf32(original);
                    b_hi[j].x[element] = high;
                    b_lo[j].x[element] = round_tf32(original - high);
                }
            }

            // Three tensor-core products retain substantially more precision
            // than a single TF32 product.
            #pragma unroll
            for (int i = 0; i < 2; ++i) {
                #pragma unroll
                for (int j = 0; j < 2; ++j) {
                    wmma::mma_sync(
                        accum[i][j], a_hi[i], b_lo[j], accum[i][j]);
                    wmma::mma_sync(
                        accum[i][j], a_lo[i], b_hi[j], accum[i][j]);
                    wmma::mma_sync(
                        accum[i][j], a_hi[i], b_hi[j], accum[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(
                storage + (warp_row + i * 16) * 64 + warp_col + j * 16,
                accum[i][j], 64, wmma::mem_row_major);
        }
    }

    __syncthreads();

    for (int index = tid; index < 4096; index += 128) {
        const int local_output_channel =
            output_channel_start + index / 64;
        const int64_t position = position_start + index % 64;

        if (local_output_channel < outputs_per_group &&
            position < output_length) {
            const int output_channel =
                group * outputs_per_group + local_output_channel;

            float value = storage[index];
            if (bias != nullptr) {
                value += bias[output_channel];
            }

            output[
                (int64_t(batch) * output_channels + output_channel)
                * output_length + position] = value;
        }
    }
#endif
}

torch::Tensor conv1d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride,
    int64_t padding,
    int64_t dilation,
    int64_t groups) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32 &&
                weight.scalar_type() == torch::kFloat32,
                "Expected float32 tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.dim() == 3 && weight.dim() == 3,
                "Expected three-dimensional tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(stride > 0 && dilation > 0 && padding >= 0 && groups > 0,
                "Invalid convolution parameters");

    const c10::cuda::CUDAGuard device_guard(x.device());
    TORCH_CHECK(at::cuda::getDeviceProperties(x.get_device())->major >= 8,
                "The custom kernel requires compute capability 8.0 or newer");

    const int64_t batch = x.size(0);
    const int64_t input_channels = x.size(1);
    const int64_t length = x.size(2);
    const int64_t output_channels = weight.size(0);
    const int64_t kernel_size = weight.size(2);

    TORCH_CHECK(input_channels == weight.size(1) * groups &&
                output_channels % groups == 0,
                "Incompatible channel dimensions");
    TORCH_CHECK(length + 2 * padding >= dilation * (kernel_size - 1) + 1,
                "Kernel is larger than the padded input");

    const int64_t output_length =
        (length + 2 * padding - dilation * (kernel_size - 1) - 1)
        / stride + 1;

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

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

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == torch::kFloat32 &&
                    bias->is_contiguous() &&
                    bias->numel() == output_channels,
                    "Invalid bias");
        bias_ptr = bias->data_ptr<float>();
    }

    TORCH_CHECK(batch * groups <= 65535, "Batch/group grid is too large");

    const dim3 grid(
        static_cast<unsigned>((output_length + 63) / 64),
        static_cast<unsigned>((output_channels / groups + 63) / 64),
        static_cast<unsigned>(batch * groups));

    conv1d_tf32_kernel<<<
        grid, 128, 0, at::cuda::getCurrentCUDAStream(x.get_device())>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        length,
        output_length,
        static_cast<int>(input_channels),
        static_cast<int>(output_channels),
        static_cast<int>(input_channels / groups),
        static_cast<int>(output_channels / groups),
        static_cast<int>(kernel_size),
        static_cast<int>(stride),
        static_cast<int>(padding),
        static_cast<int>(dilation),
        static_cast<int>(groups));

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="custom_conv1d_implicit_tf32_v1",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["conv1d_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.ndim == 3
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and not isinstance(conv.padding, str)
            and not torch.is_autocast_enabled()
            and x.shape[0] * conv.groups <= 65535
            and torch.cuda.get_device_capability(x.device)[0] >= 8
        )

        if not supported:
            return conv(x)

        return _get_extension().conv1d_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias,
            conv.stride[0],
            conv.padding[0],
            conv.dilation[0],
            conv.groups,
        )
```