```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,
    int64_t stride,
    int64_t padding,
    int64_t dilation,
    int64_t groups);
"""


_cuda_source = r"""
#include <torch/extension.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <mma.h>

namespace {

constexpr int TM = 32;
constexpr int TN = 128;
constexpr int TK = 32;
constexpr int AS = TK + 8;
constexpr int BS = TN + 8;

union SharedStorage {
    struct {
        half ah[TM * AS];
        half al[TM * AS];
        half bh[TK * BS];
        half bl[TK * BS];
    } operands;
    float output[TM * TN];
};

__global__ void conv1d_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int input_channels,
    int output_channels,
    int input_length,
    int output_length,
    int channels_per_group,
    int outputs_per_group,
    int kernel_size,
    int stride,
    int padding,
    int dilation,
    int groups)
{
    const int tid = threadIdx.x;
    const int batch = blockIdx.z / groups;
    const int group = blockIdx.z % groups;
    const int channel_base = blockIdx.y * TM;
    const int position_base = blockIdx.x * TN;
    const int reduction_size = channels_per_group * kernel_size;

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
    using namespace nvcuda;

    __shared__ __align__(32) SharedStorage storage;

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

    wmma::fragment<wmma::accumulator, 16, 16, 16, float> main0;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> main1;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> correction0;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> correction1;

    wmma::fill_fragment(main0, 0.0f);
    wmma::fill_fragment(main1, 0.0f);
    wmma::fill_fragment(correction0, 0.0f);
    wmma::fill_fragment(correction1, 0.0f);

    for (int reduction_base = 0;
         reduction_base < reduction_size;
         reduction_base += TK)
    {
        for (int index = tid; index < TM * TK; index += blockDim.x) {
            const int row = index / TK;
            const int k = index % TK;
            const int channel = channel_base + row;
            const int reduction = reduction_base + k;

            float value = 0.0f;
            if (channel < outputs_per_group && reduction < reduction_size) {
                const int global_channel =
                    group * outputs_per_group + channel;
                value = weight[
                    static_cast<int64_t>(global_channel) * reduction_size
                    + reduction];
            }

            const half high = __float2half_rn(value);
            storage.operands.ah[row * AS + k] = high;
            storage.operands.al[row * AS + k] =
                __float2half_rn(value - __half2float(high));
        }

        for (int index = tid; index < TK * TN; index += blockDim.x) {
            const int k = index / TN;
            const int col = index % TN;
            const int reduction = reduction_base + k;
            const int position = position_base + col;

            float value = 0.0f;
            if (reduction < reduction_size && position < output_length) {
                const int input_channel = reduction / kernel_size;
                const int kernel_position = reduction % kernel_size;
                const int64_t input_position =
                    static_cast<int64_t>(position) * stride - padding
                    + static_cast<int64_t>(kernel_position) * dilation;

                if (input_position >= 0 && input_position < input_length) {
                    const int global_input_channel =
                        group * channels_per_group + input_channel;
                    value = x[
                        (static_cast<int64_t>(batch) * input_channels
                         + global_input_channel) * input_length
                        + input_position];
                }
            }

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

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < TK; k += 16) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 16, half, wmma::row_major> ah;
            wmma::fragment<
                wmma::matrix_a, 16, 16, 16, half, wmma::row_major> al;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 16, half, wmma::row_major> bh0;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 16, half, wmma::row_major> bh1;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 16, half, wmma::row_major> bl0;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 16, half, wmma::row_major> bl1;

            wmma::load_matrix_sync(
                ah, storage.operands.ah + warp_row * AS + k, AS);
            wmma::load_matrix_sync(
                al, storage.operands.al + warp_row * AS + k, AS);
            wmma::load_matrix_sync(
                bh0, storage.operands.bh + k * BS + warp_col, BS);
            wmma::load_matrix_sync(
                bh1, storage.operands.bh + k * BS + warp_col + 16, BS);
            wmma::load_matrix_sync(
                bl0, storage.operands.bl + k * BS + warp_col, BS);
            wmma::load_matrix_sync(
                bl1, storage.operands.bl + k * BS + warp_col + 16, BS);

            wmma::mma_sync(main0, ah, bh0, main0);
            wmma::mma_sync(main1, ah, bh1, main1);
            wmma::mma_sync(correction0, ah, bl0, correction0);
            wmma::mma_sync(correction1, ah, bl1, correction1);
            wmma::mma_sync(correction0, al, bh0, correction0);
            wmma::mma_sync(correction1, al, bh1, correction1);
        }

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < main0.num_elements; ++i) {
        main0.x[i] += correction0.x[i];
        main1.x[i] += correction1.x[i];
    }

    wmma::store_matrix_sync(
        storage.output + warp_row * TN + warp_col,
        main0, TN, wmma::mem_row_major);
    wmma::store_matrix_sync(
        storage.output + warp_row * TN + warp_col + 16,
        main1, TN, wmma::mem_row_major);

    __syncthreads();

    for (int index = tid; index < TM * TN; index += blockDim.x) {
        const int channel = channel_base + index / TN;
        const int position = position_base + index % TN;

        if (channel < outputs_per_group && position < output_length) {
            const int global_channel = group * outputs_per_group + channel;
            float value = storage.output[index];
            if (bias != nullptr) {
                value += bias[global_channel];
            }
            output[
                (static_cast<int64_t>(batch) * output_channels
                 + global_channel) * output_length + position] = value;
        }
    }
#else
    for (int index = tid; index < TM * TN; index += blockDim.x) {
        const int channel = channel_base + index / TN;
        const int position = position_base + index % TN;

        if (channel < outputs_per_group && position < output_length) {
            const int global_channel = group * outputs_per_group + channel;
            float value = 0.0f;

            for (int ic = 0; ic < channels_per_group; ++ic) {
                for (int k = 0; k < kernel_size; ++k) {
                    const int64_t input_position =
                        static_cast<int64_t>(position) * stride
                        - padding + static_cast<int64_t>(k) * dilation;

                    if (input_position >= 0 && input_position < input_length) {
                        const int64_t input_index =
                            (static_cast<int64_t>(batch) * input_channels
                             + group * channels_per_group + ic)
                            * input_length + input_position;
                        const int64_t weight_index =
                            static_cast<int64_t>(global_channel)
                            * reduction_size + ic * kernel_size + k;
                        value = fmaf(x[input_index], weight[weight_index], value);
                    }
                }
            }

            if (bias != nullptr) {
                value += bias[global_channel];
            }
            output[
                (static_cast<int64_t>(batch) * output_channels
                 + global_channel) * output_length + position] = value;
        }
    }
#endif
}

} // namespace

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.dim() == 3 && weight.dim() == 3,
                "Expected three-dimensional tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(stride > 0 && dilation > 0 && padding >= 0 && groups > 0,
                "Invalid convolution parameters");
    TORCH_CHECK(x.size(1) == weight.size(1) * groups &&
                weight.size(0) % groups == 0,
                "Incompatible convolution dimensions");

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t effective_kernel = dilation * (weight.size(2) - 1) + 1;
    TORCH_CHECK(x.size(2) + 2 * padding >= effective_kernel,
                "Kernel exceeds padded input length");

    const int64_t output_length =
        (x.size(2) + 2 * padding - effective_kernel) / stride + 1;

    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() == weight.size(0),
                    "Invalid bias");
        bias_ptr = bias->data_ptr<float>();
    }

    auto output = torch::empty(
        {x.size(0), weight.size(0), output_length}, x.options());

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

    const int64_t outputs_per_group = weight.size(0) / groups;
    TORCH_CHECK(x.size(0) * groups <= 65535 &&
                (outputs_per_group + TM - 1) / TM <= 65535,
                "Convolution exceeds CUDA grid limits");

    const dim3 grid(
        (output_length + TN - 1) / TN,
        (outputs_per_group + TM - 1) / TM,
        x.size(0) * groups);

    conv1d_kernel<<<grid, 256, 0, c10::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        x.size(1),
        weight.size(0),
        x.size(2),
        output_length,
        weight.size(1),
        outputs_per_group,
        weight.size(2),
        stride,
        padding,
        dilation,
        groups);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""


_conv1d_extension = load_inline(
    name="implicit_conv1d_compensated_wmma",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv1d_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _Conv1dFunction(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias, stride, padding, dilation, groups):
        ctx.save_for_backward(x, weight)
        ctx.parameters = (stride, padding, dilation, groups)
        ctx.has_bias = bias is not None
        return _conv1d_extension.conv1d_cuda(
            x, weight, bias, stride, padding, dilation, groups
        )

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        stride, padding, dilation, groups = ctx.parameters
        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,
                stride, padding, dilation, groups
            )
        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv1d_weight(
                x, weight.shape, grad_output,
                stride, padding, dilation, groups
            )
        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, None, None, None, None


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
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.ndim != 3
            or isinstance(conv.padding, str)
            or torch.is_autocast_enabled()
            or x.shape[0] * conv.groups > 65535
            or (conv.out_channels // conv.groups + 31) // 32 > 65535
        ):
            return conv(x)

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