def gen_forward():

    kernels = [3, 5, 7, 15, 31, 63, 127, 255]
    blocks = [32, 64, 128, 256]

    head = """
/**
 * Copyright (c) Facebook, Inc. and its affiliates.
 *
 * This source code is licensed under the MIT license found in the
 * LICENSE file in the root directory of this source tree.
 */

#include "dynamicconv_cuda.cuh"

std::vector<at::Tensor> dynamicconv_cuda_forward(at::Tensor input, at::Tensor weight, int padding_l) {

    at::DeviceGuard g(input.device());
    const auto minibatch = input.size(0);
    const auto numFeatures = input.size(1);
    const auto sequenceLength = input.size(2);

    const auto numHeads = weight.size(1);
    const auto filterSize = weight.size(2);

    const auto numFiltersInBlock = numFeatures / numHeads;
    const dim3 blocks(minibatch, numFeatures);

    auto output = at::zeros_like(input);
    auto stream = at::cuda::getCurrentCUDAStream();
"""

    switch = """
    switch(filterSize) {
"""

    case_k = """
        case {k}:
"""

    main_block = """
            if (padding_l == {pad}) {{
                AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), "dynamicconv_forward", ([&] {{
                    dynamicconv_forward_kernel<{k}, {b_size}, {pad}, scalar_t>
                    <<<blocks, {b_size}, 0, stream>>>(
                            input.data<scalar_t>(),
                            weight.data<scalar_t>(),
                            minibatch,
                            sequenceLength,
                            numFeatures,
                            numFiltersInBlock,
                            numHeads,
                            output.data<scalar_t>());
                }}));
            }} else
"""

    bad_padding = """
            {
                std::cout << "WARNING: Unsupported padding size - skipping forward pass" << std::endl;
            }
            break;\n
"""

    end = """
        default:
            std::cout << "WARNING: Unsupported filter length passed - skipping forward pass" << std::endl;
    }

    return {output};
}
"""

    with open("dynamicconv_cuda_forward.cu", "w") as forward:
        forward.write(head)
        forward.write(switch)
        for k in kernels:
            b_size = 32
            for b in blocks:
                if b > k:
                    b_size = b
                    break
            forward.write(case_k.format(k=k))
            for pad in [k // 2, k - 1]:
                forward.write(main_block.format(k=k, b_size=b_size, pad=pad))
            forward.write(bad_padding)
        forward.write(end)
switch =
fairseq/modules/dynamicconv_layer/cuda_function_gen.py