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

constexpr int TILE_M = 64;
constexpr int TILE_N = 128;
constexpr int TILE_K = 16;

template<int FIXED_KERNEL>
__global__ __launch_bounds__(256)
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 channel_tiles)
{
    __shared__ float ws[TILE_M][TILE_K + 1];
    __shared__ float xs[TILE_K][TILE_N];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int row_lane = tid >> 5;

    const int batch = blockIdx.z;
    const int group = blockIdx.y / channel_tiles;
    const int channel_tile = blockIdx.y - group * channel_tiles;
    const int local_channel_start = channel_tile * TILE_M;
    const int channel_start =
        group * outputs_per_group + local_channel_start;
    const int position_start = blockIdx.x * TILE_N;

    const int ks = FIXED_KERNEL ? FIXED_KERNEL : kernel_size;
    const int reduction_size = channels_per_group * ks;

    const float* group_input =
        x + ((int64_t)batch * input_channels
             + group * channels_per_group) * input_length;

    float acc[8][4];

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

    for (int base_k = 0; base_k < reduction_size; base_k += TILE_K) {
        #pragma unroll
        for (int offset = 0; offset < TILE_M * TILE_K; offset += 256) {
            const int index = tid + offset;
            const int row = index / TILE_K;
            const int kk = index % TILE_K;
            const int k = base_k + kk;

            float value = 0.0f;
            if (local_channel_start + row < outputs_per_group
                && k < reduction_size) {
                value = weight[
                    (int64_t)(channel_start + row) * reduction_size + k];
            }
            ws[row][kk] = value;
        }

        #pragma unroll
        for (int offset = 0; offset < TILE_K * TILE_N; offset += 256) {
            const int index = tid + offset;
            const int kk = index / TILE_N;
            const int column = index % TILE_N;
            const int k = base_k + kk;
            const int output_position = position_start + column;

            float value = 0.0f;
            if (k < reduction_size && output_position < output_length) {
                const int input_channel = k / ks;
                const int tap = k - input_channel * ks;
                const int64_t input_position =
                    (int64_t)output_position * stride
                    - padding + (int64_t)tap * dilation;

                if (input_position >= 0 && input_position < input_length) {
                    value = group_input[
                        (int64_t)input_channel * input_length + input_position];
                }
            }
            xs[kk][column] = value;
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < TILE_K; ++kk) {
            const float b0 = xs[kk][lane];
            const float b1 = xs[kk][lane + 32];
            const float b2 = xs[kk][lane + 64];
            const float b3 = xs[kk][lane + 96];

            #pragma unroll
            for (int i = 0; i < 8; ++i) {
                const float a = ws[row_lane + i * 8][kk];
                acc[i][0] = fmaf(a, b0, acc[i][0]);
                acc[i][1] = fmaf(a, b1, acc[i][1]);
                acc[i][2] = fmaf(a, b2, acc[i][2]);
                acc[i][3] = fmaf(a, b3, acc[i][3]);
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        const int local_channel = local_channel_start + row_lane + i * 8;
        const int output_channel = channel_start + row_lane + i * 8;

        if (local_channel < outputs_per_group) {
            const float bias_value = bias ? bias[output_channel] : 0.0f;
            const int64_t output_base =
                ((int64_t)batch * output_channels + output_channel)
                * output_length;

            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                const int position = position_start + lane + j * 32;
                if (position < output_length) {
                    output[output_base + position] =
                        acc[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(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat
                && weight.scalar_type() == at::kFloat,
                "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");

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

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

    TORCH_CHECK(input_channels == channels_per_group * groups,
                "Input channel mismatch");
    TORCH_CHECK(output_channels % groups == 0, "Invalid output groups");
    TORCH_CHECK(kernel_size > 0, "Kernel size must be positive");

    const int64_t numerator =
        input_length + 2 * padding - dilation * (kernel_size - 1) - 1;
    TORCH_CHECK(numerator >= 0, "Kernel exceeds padded input length");
    const int64_t output_length = numerator / stride + 1;
    const int64_t outputs_per_group = output_channels / groups;

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

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

    const int channel_tiles = (outputs_per_group + TILE_M - 1) / TILE_M;
    TORCH_CHECK(batch <= 65535 && groups * channel_tiles <= 65535,
                "Convolution exceeds CUDA grid limits");

    const dim3 grid(
        (output_length + TILE_N - 1) / TILE_N,
        groups * channel_tiles,
        batch);
    const auto stream = at::cuda::getCurrentCUDAStream();

    if (kernel_size == 3) {
        conv1d_kernel<3><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            input_channels, output_channels, input_length, output_length,
            channels_per_group, outputs_per_group, kernel_size,
            stride, padding, dilation, channel_tiles);
    } else {
        conv1d_kernel<0><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            input_channels, output_channels, input_length, output_length,
            channels_per_group, outputs_per_group, kernel_size,
            stride, padding, dilation, channel_tiles);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_conv1d_extension = load_inline(
    name="tiled_custom_conv1d_fp32",
    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.dim() != 3
            or isinstance(conv.padding, str)
            or torch.is_autocast_enabled()
        ):
            return conv(x)

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