```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

# Custom implicit-GEMM Conv1d (valid, stride=1) using TF32 tensor cores (WMMA)
# with FP32 accumulation. Specialized for C_out % 128 == 0, (C_in*K) % 32 == 0.
_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <mma.h>

using namespace nvcuda;

#define BM 128
#define BN 256
#define BK 32
#define BS_LD 264   // padded leading dimension of staged B tile (BK x BS_LD floats)

// C[o, n] = sum_kd A[o, kd] * B[kd, n]
//   A[o, kd]        = w[o, kd]                    (kd = c*K + k)
//   B[kd, n]        = x[b, c, l + k], l = n0+nn
// out[b, o, l]      = conv1d(x, w)[b, o, l]
__global__ void __launch_bounds__(512, 1)
conv1d_wmma_kernel(const float* __restrict__ x,
                   const float* __restrict__ w,
                   float* __restrict__ out,
                   int B, int C_in, int L, int L_out, int L_pad,
                   int C_out, int K, int K_dim) {
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
    __shared__ float Bs[BK][BS_LD];

    const int n0 = blockIdx.x * BN;
    const int m0 = blockIdx.y * BM;
    const int b  = blockIdx.z;

    const int tid  = threadIdx.x;
    const int warp = tid >> 5;
    const int wm   = warp >> 2;   // 0..3 (rows)
    const int wn   = warp & 3;    // 0..3 (cols)

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[2][4];
    #pragma unroll
    for (int i = 0; i < 2; ++i) {
        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            wmma::fill_fragment(acc[i][j], 0.0f);
        }
    }

    const float* __restrict__ xbase = x + (size_t)b * (size_t)C_in * (size_t)L;
    const int nchunks = K_dim / BK;   // exact (guaranteed by host checks)

    for (int ch = 0; ch < nchunks; ++ch) {
        const int kd0 = ch * BK;

        // Cooperative load of the B tile (BK x BN) with implicit im2col gather.
        #pragma unroll
        for (int it = 0; it < (BK * BN) / 512; ++it) {
            const int idx = tid + it * 512;
            const int kk  = idx >> 8;    // 0..31
            const int nn  = idx & 255;   // 0..255
            const int kd  = kd0 + kk;
            int c, k;
            if (K == 3) {
                c = (int)(__umulhi((unsigned)kd, 0xAAAAAAABu) >> 1);
                k = kd - c * 3;
            } else {
                c = kd / K;
                k = kd - c * K;
            }
            const int l = n0 + nn;
            float v = 0.0f;
            if (l < L_out) {
                v = xbase[(size_t)c * (size_t)L + (size_t)l + k];
            }
            Bs[kk][nn] = v;
        }
        __syncthreads();

        const int o0 = m0 + wm * 32;
        #pragma unroll
        for (int ks = 0; ks < BK / 8; ++ks) {
            const int kd = kd0 + ks * 8;
            wmma::fragment<wmma::matrix_a, 16, 16, 8, wmma::precision::tf32, wmma::row_major> af[2];
            wmma::load_matrix_sync(af[0], w + (size_t)o0 * (size_t)K_dim + kd, K_dim);
            wmma::load_matrix_sync(af[1], w + (size_t)(o0 + 16) * (size_t)K_dim + kd, K_dim);
            #pragma unroll
            for (int e = 0; e < af[0].num_elements; ++e) {
                unsigned ua = __float_as_uint(af[0].x[e]) & 0xFFFFE000u;
                unsigned ub = __float_as_uint(af[1].x[e]) & 0xFFFFE000u;
                af[0].x[e] = __uint_as_float(ua);
                af[1].x[e] = __uint_as_float(ub);
            }
            #pragma unroll
            for (int n = 0; n < 4; ++n) {
                wmma::fragment<wmma::matrix_b, 16, 16, 8, wmma::precision::tf32, wmma::row_major> bf;
                wmma::load_matrix_sync(bf, &Bs[ks * 8][wn * 64 + n * 16], BS_LD);
                #pragma unroll
                for (int e = 0; e < bf.num_elements; ++e) {
                    unsigned u = __float_as_uint(bf.x[e]) & 0xFFFFE000u;
                    bf.x[e] = __uint_as_float(u);
                }
                wmma::mma_sync(acc[0][n], af[0], bf, acc[0][n]);
                wmma::mma_sync(acc[1][n], af[1], bf, acc[1][n]);
            }
        }
        __syncthreads();
    }

    // Epilogue: padded output buffer guarantees full 16x16 tiles and a
    // store leading dimension that is a multiple of 8.
    float* outbase = out + ((size_t)b * (size_t)C_out + (size_t)(m0 + wm * 32)) * (size_t)L_pad
                     + (size_t)n0 + wn * 64;
    #pragma unroll
    for (int m = 0; m < 2; ++m) {
        #pragma unroll
        for (int n = 0; n < 4; ++n) {
            wmma::store_matrix_sync(outbase + (size_t)m * 16 * (size_t)L_pad + n * 16,
                                    acc[m][n], (unsigned)L_pad, wmma::mem_row_major);
        }
    }
#endif
}

torch::Tensor conv1d_forward(torch::Tensor x, torch::Tensor w) {
    TORCH_CHECK(x.is_cuda() && w.is_cuda(), "inputs must be CUDA");
    TORCH_CHECK(x.scalar_type() == at::kFloat && w.scalar_type() == at::kFloat,
                "inputs must be float32");
    TORCH_CHECK(x.dim() == 3 && w.dim() == 3, "bad ranks");
    const int B     = (int)x.size(0);
    const int C_in  = (int)x.size(1);
    const int L     = (int)x.size(2);
    const int C_out = (int)w.size(0);
    const int K     = (int)w.size(2);
    const int K_dim = C_in * K;
    const int L_out = L - K + 1;
    TORCH_CHECK(L_out > 0, "kernel larger than length");
    TORCH_CHECK((K_dim % BK) == 0, "C_in*K must be divisible by 32");
    TORCH_CHECK((C_out % BM) == 0, "C_out must be divisible by 128");
    TORCH_CHECK(B <= 65535, "batch too large");

    const int L_pad = ((L_out + BN - 1) / BN) * BN;
    auto out = at::empty({B, C_out, (long)L_pad}, x.options());

    dim3 grid((L_out + BN - 1) / BN, C_out / BM, B);
    auto stream = at::cuda::getCurrentCUDAStream();
    conv1d_wmma_kernel<<<grid, 512, 0, stream>>>(
        x.data_ptr<float>(), w.data_ptr<float>(), out.data_ptr<float>(),
        B, C_in, L, L_out, L_pad, C_out, K, K_dim);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    // Return a view with the exact expected output shape.
    return out.narrow(2, 0, L_out);
}
"""

_cpp_src = "torch::Tensor conv1d_forward(torch::Tensor x, torch::Tensor w);"

_conv_ext = load_inline(
    name="conv1d_wmma_ext",
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=["conv1d_forward"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class Model(nn.Module):
    """
    Optimized 1D convolution. Uses a custom implicit-GEMM CUDA kernel with
    TF32 tensor cores (FP32 accumulation) when the problem configuration is
    supported; falls back to the standard nn.Conv1d otherwise.
    """
    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(Model, self).__init__()
        self.conv1d = nn.Conv1d(in_channels, out_channels, kernel_size,
                                stride=stride, padding=padding, dilation=dilation,
                                groups=groups, bias=bias)
        K = kernel_size[0] if isinstance(kernel_size, (tuple, list)) else kernel_size

        ok = (stride == 1 and padding == 0 and dilation == 1 and groups == 1
              and (not bias)
              and (in_channels * K) % 32 == 0
              and out_channels % 128 == 0)
        if ok:
            try:
                ok = (torch.cuda.is_available()
                      and torch.cuda.get_device_capability(0)[0] >= 8)
            except Exception:
                ok = False
        self._use_custom = ok

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if (self._use_custom and x.is_cuda and x.dtype == torch.float32
                and x.dim() == 3 and x.size(0) <= 65535
                and x.size(2) >= self.conv1d.weight.size(2)):
            return _conv_ext.conv1d_forward(x.contiguous(),
                                            self.conv1d.weight.contiguous())
        return self.conv1d(x)


# Alias required by the task.
class ModelNew(Model):
    pass
```