```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 <algorithm>

template<bool SPECIALIZED>
__global__ void conv1d_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int64_t length,
    int64_t output_length,
    int channels,
    int output_channels,
    int channels_per_group,
    int outputs_per_group,
    int kernel_size,
    int stride,
    int padding,
    int dilation,
    int groups)
{
    constexpr int TILE_L = 128;
    constexpr int TILE_C = 64;
    constexpr int TILE_K = 16;

    __shared__ float xs[TILE_K][TILE_L];
    __shared__ float ws[TILE_K][TILE_C + 1];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int tid = ty * 16 + tx;

    const int group = blockIdx.z % groups;
    const int batch = blockIdx.z / groups;
    const int first_oc = blockIdx.y * TILE_C;
    const int64_t first_l = int64_t(blockIdx.x) * TILE_L;
    const int reduction_size = channels_per_group * kernel_size;

    const int64_t input_base =
        (int64_t(batch) * channels + group * channels_per_group) * length;

    float accum[4][8];

    #pragma unroll
    for (int i = 0; i < 4; ++i) {
        #pragma unroll
        for (int j = 0; j < 8; ++j) {
            accum[i][j] = 0.0f;
        }
    }

    for (int base_k = 0; base_k < reduction_size; base_k += TILE_K) {
        #pragma unroll
        for (int index = tid; index < TILE_K * TILE_L; index += 256) {
            const int k_local = index / TILE_L;
            const int l_local = index % TILE_L;
            const int k = base_k + k_local;
            const int64_t out_l = first_l + l_local;

            int ic;
            int offset;
            int64_t in_l;

            if (SPECIALIZED) {
                ic = k / 3;
                offset = k % 3;
                in_l = out_l + offset;
            } else {
                ic = k / kernel_size;
                offset = k % kernel_size;
                in_l = out_l * stride - padding
                     + int64_t(offset) * dilation;
            }

            float value = 0.0f;
            if (k < reduction_size &&
                out_l < output_length &&
                in_l >= 0 && in_l < length) {
                value = x[input_base + int64_t(ic) * length + in_l];
            }
            xs[k_local][l_local] = value;
        }

        #pragma unroll
        for (int index = tid; index < TILE_C * TILE_K; index += 256) {
            const int oc_local = index / TILE_K;
            const int k_local = index % TILE_K;
            const int oc = first_oc + oc_local;
            const int k = base_k + k_local;

            float value = 0.0f;
            if (oc < outputs_per_group && k < reduction_size) {
                const int global_oc = group * outputs_per_group + oc;
                value = weight[int64_t(global_oc) * reduction_size + k];
            }
            ws[k_local][oc_local] = value;
        }

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < TILE_K; ++k) {
            float xv[8];
            float wv[4];

            #pragma unroll
            for (int j = 0; j < 8; ++j) {
                xv[j] = xs[k][tx + j * 16];
            }

            #pragma unroll
            for (int i = 0; i < 4; ++i) {
                wv[i] = ws[k][ty + i * 16];
            }

            #pragma unroll
            for (int i = 0; i < 4; ++i) {
                #pragma unroll
                for (int j = 0; j < 8; ++j) {
                    accum[i][j] = fmaf(wv[i], xv[j], accum[i][j]);
                }
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 4; ++i) {
        const int oc = first_oc + ty + i * 16;
        if (oc < outputs_per_group) {
            const int global_oc = group * outputs_per_group + oc;
            const float b = bias ? bias[global_oc] : 0.0f;
            const int64_t output_base =
                (int64_t(batch) * output_channels + global_oc) * output_length;

            #pragma unroll
            for (int j = 0; j < 8; ++j) {
                const int64_t out_l = first_l + tx + j * 16;
                if (out_l < output_length) {
                    output[output_base + out_l] = accum[i][j] + b;
                }
            }
        }
    }
}

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.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32 &&
                weight.scalar_type() == torch::kFloat32,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 3 && weight.dim() == 3, "Expected rank-three tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(stride > 0 && dilation > 0 && padding >= 0 && groups > 0,
                "Invalid convolution parameters");

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

    const int64_t batch = x.size(0);
    const int64_t 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(channels == weight.size(1) * groups &&
                output_channels % groups == 0,
                "Invalid grouped convolution dimensions");

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

    const int64_t output_length =
        (length + 2 * padding - effective_kernel) / 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 count exceeds grid limit");

    const dim3 block(16, 16);
    const dim3 grid(
        (output_length + 127) / 128,
        (output_channels / groups + 63) / 64,
        batch * groups);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (kernel_size == 3 && stride == 1 && padding == 0 && dilation == 1) {
        conv1d_kernel<true><<<grid, block, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(), length, output_length,
            channels, output_channels, channels / groups,
            output_channels / groups, kernel_size, stride, padding,
            dilation, groups);
    } else {
        conv1d_kernel<false><<<grid, block, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(), length, output_length,
            channels, output_channels, channels / groups,
            output_channels / groups, kernel_size, stride, padding,
            dilation, groups);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="tiled_implicit_conv1d_ext",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["conv1d_cuda"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _extension


class _Conv1dFunction(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias, stride, padding, dilation, groups):
        ctx.save_for_backward(x, weight)
        ctx.params = (stride, padding, dilation, groups)
        ctx.has_bias = bias is not None
        return _get_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.params
        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
        ):
            return conv(x)

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