```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 <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <climits>

constexpr int BM = 64;
constexpr int BN = 128;
constexpr int BK = 16;

template<int STATIC_KERNEL>
__global__ __launch_bounds__(256)
void conv1d_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int in_channels,
    int out_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,
    int spatial_tiles,
    int channel_tiles)
{
    __shared__ float tile_a[BK][BM];
    __shared__ float tile_b[BK][BN];

    const int tid = threadIdx.x;
    const int tx = tid & 15;
    const int ty = tid >> 4;

    int64_t tile = blockIdx.x;
    const int spatial_tile = tile % spatial_tiles;
    tile /= spatial_tiles;
    const int channel_tile = tile % channel_tiles;
    tile /= channel_tiles;
    const int group = tile % groups;
    const int batch = tile / groups;

    const int row_base = channel_tile * BM;
    const int col_base = spatial_tile * BN;
    const int ks = STATIC_KERNEL ? STATIC_KERNEL : kernel_size;
    const int reduction_size = channels_per_group * ks;

    const float* input_base =
        x + (static_cast<int64_t>(batch) * in_channels
             + group * channels_per_group) * input_length;
    const float* weight_base =
        weight + static_cast<int64_t>(group) *
                 outputs_per_group * reduction_size;

    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 k_base = 0; k_base < reduction_size; k_base += BK) {
        #pragma unroll
        for (int index = tid; index < BM * BK; index += 256) {
            const int row = index / BK;
            const int k = index % BK;
            const int global_row = row_base + row;
            const int global_k = k_base + k;

            float value = 0.0f;
            if (global_row < outputs_per_group &&
                global_k < reduction_size) {
                value = weight_base[
                    static_cast<int64_t>(global_row) * reduction_size
                    + global_k];
            }
            tile_a[k][row] = value;
        }

        #pragma unroll
        for (int index = tid; index < BK * BN; index += 256) {
            const int k = index / BN;
            const int col = index % BN;
            const int global_k = k_base + k;
            const int global_col = col_base + col;

            float value = 0.0f;
            if (global_k < reduction_size && global_col < output_length) {
                const int channel = global_k / ks;
                const int tap = global_k % ks;
                const int64_t position =
                    static_cast<int64_t>(global_col) * stride
                    - padding + static_cast<int64_t>(tap) * dilation;

                if (position >= 0 && position < input_length) {
                    value = input_base[
                        static_cast<int64_t>(channel) * input_length
                        + position];
                }
            }
            tile_b[k][col] = value;
        }

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < BK; ++k) {
            float a[4];
            float b[8];

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

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

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

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 4; ++i) {
        const int local_channel = row_base + ty + i * 16;
        if (local_channel < outputs_per_group) {
            const int channel = group * outputs_per_group + local_channel;
            const float bias_value = bias ? bias[channel] : 0.0f;
            const int64_t output_base =
                (static_cast<int64_t>(batch) * out_channels + channel)
                * output_length;

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

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(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Custom convolution requires float32 tensors");
    TORCH_CHECK(x.dim() == 3 && weight.dim() == 3,
                "Expected three-dimensional input and weight");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(groups > 0 && stride > 0 && dilation > 0 && padding >= 0,
                "Invalid convolution parameters");

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

    x = x.contiguous();
    weight = weight.contiguous();

    const int64_t batch = x.size(0);
    const int64_t in_channels = x.size(1);
    const int64_t input_length = x.size(2);
    const int64_t out_channels = weight.size(0);
    const int64_t channels_per_group = weight.size(1);
    const int64_t kernel_size = weight.size(2);

    TORCH_CHECK(in_channels == channels_per_group * groups &&
                out_channels % groups == 0,
                "Invalid channel or group dimensions");
    TORCH_CHECK(channels_per_group > 0 && out_channels > 0 &&
                kernel_size > 0 && input_length > 0,
                "Invalid convolution dimensions");

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

    const int64_t output_length =
        (input_length + 2 * padding - effective_kernel) / stride + 1;
    const int64_t outputs_per_group = out_channels / groups;

    TORCH_CHECK(in_channels <= INT_MAX && out_channels <= INT_MAX &&
                input_length <= INT_MAX && output_length <= INT_MAX &&
                channels_per_group * kernel_size <= INT_MAX &&
                stride <= INT_MAX && padding <= INT_MAX &&
                dilation <= INT_MAX && groups <= INT_MAX,
                "Convolution dimensions exceed supported indexing limits");

    torch::Tensor contiguous_bias;
    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == at::kFloat &&
                    bias->dim() == 1 && bias->numel() == out_channels,
                    "Invalid bias");
        contiguous_bias = bias->contiguous();
        bias_ptr = contiguous_bias.data_ptr<float>();
    }

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

    if (batch == 0) {
        return output;
    }

    const int64_t spatial_tiles = (output_length + BN - 1) / BN;
    const int64_t channel_tiles = (outputs_per_group + BM - 1) / BM;
    const int64_t blocks =
        batch * groups * channel_tiles * spatial_tiles;

    TORCH_CHECK(blocks <= INT_MAX, "Convolution grid is too large");

    const cudaStream_t stream =
        at::cuda::getCurrentCUDAStream(x.get_device());

    if (kernel_size == 3) {
        conv1d_kernel<3><<<static_cast<unsigned int>(blocks), 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            in_channels, out_channels, input_length, output_length,
            channels_per_group, outputs_per_group, kernel_size,
            stride, padding, dilation, groups, spatial_tiles, channel_tiles);
    } else {
        conv1d_kernel<0><<<static_cast<unsigned int>(blocks), 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            in_channels, out_channels, input_length, output_length,
            channels_per_group, outputs_per_group, kernel_size,
            stride, padding, dilation, groups, spatial_tiles, channel_tiles);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_conv1d_extension = load_inline(
    name="tiled_conv1d_cuda_v1",
    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.stride = stride
        ctx.padding = padding
        ctx.dilation = dilation
        ctx.groups = 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
        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=ctx.stride,
                padding=ctx.padding,
                dilation=ctx.dilation,
                groups=ctx.groups,
            )

        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv1d_weight(
                x,
                weight.shape,
                grad_output,
                stride=ctx.stride,
                padding=ctx.padding,
                dilation=ctx.dilation,
                groups=ctx.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()
        ):
            return conv(x)

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