/*
 * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved.
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 *     http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

#pragma once
#include "fused_multihead_attention.h"
#include "fused_multihead_attention_common.h"
#include <assert.h>
#include <stdint.h>

namespace fastertransformer {
struct Fused_multihead_attention_params_v2 {
    // The QKV matrices.
    void* qkv_ptr;
    // The mask to implement drop-out.
    void* packed_mask_ptr;

    // The O matrix (output).
    void* o_ptr;

    // The stride between rows of the Q, K and V matrices.
    int64_t qkv_stride_in_bytes;
    // The stride between matrices of packed mask.
    int64_t packed_mask_stride_in_bytes;
    // The stride between rows of O.
    int64_t o_stride_in_bytes;

#if defined(STORE_P)
    // The pointer to the P matrix (for debugging).
    void* p_ptr;
    // The stride between rows of the P matrix (for debugging).
    int64_t p_stride_in_bytes;
#endif  // defined(STORE_P)

#if defined(STORE_S)
    // The pointer to the S matrix (for debugging).
    void* s_ptr;
    // The stride between rows of the S matrix (for debugging).
    int64_t s_stride_in_bytes;
#endif  // defined(STORE_S)

    // The dimensions.
    int b, h, s, d;
    // The scaling factors for the kernel.
    uint32_t scale_bmm1, scale_softmax, scale_bmm2;

    // Do we use Niall's trick to avoid I2F/F2I in the INT8 kernel.
    // See https://confluence.nvidia.com/pages/viewpage.action?pageId=302779721 for details.
    bool enable_i2f_trick;

    // array of length b+1 holding prefix sum of actual sequence lenghts
    int* cu_seqlens;

    // use C/32 Format.
    bool interleaved        = false;
    bool ignore_b1opt       = false;
    bool force_unroll       = false;
    bool use_int8_scale_max = false;

    //***The additional paramters for fused_mha_with_relPosBias kernels
    // The relative position bias.
    void* packed_relative_position_bias_ptr = nullptr;

    int window_num, actual_seqlen;

    void clear()
    {
        qkv_ptr                           = nullptr;
        packed_mask_ptr                   = nullptr;
        packed_relative_position_bias_ptr = nullptr;
        o_ptr                             = nullptr;

        qkv_stride_in_bytes         = 0;
        packed_mask_stride_in_bytes = 0;
        o_stride_in_bytes           = 0;
#if defined(STORE_P)
        p_ptr             = nullptr;
        p_stride_in_bytes = 0
#endif  // defined(STORE_P)

#if defined(STORE_S)
            s_ptr         = nullptr;
        s_stride_in_bytes = 0;
#endif  // defined(STORE_S)

        b = 0;
        h = 0;
        s = 0;
        d = 0;
        // The scaling factors for the kernel.
        scale_bmm1    = 0;
        scale_softmax = 0;
        scale_bmm2    = 0;

        enable_i2f_trick = false;

        cu_seqlens         = nullptr;
        interleaved        = false;
        ignore_b1opt       = false;
        force_unroll       = false;
        use_int8_scale_max = false;
    }
};
////////////////////////////////////////////////////////////////////////////////////////////////////
extern unsigned char fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin[];
extern unsigned char fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_32_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_64_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_32_sm75_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_64_sm75_cu_cubin[];

extern unsigned char fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin[];

extern unsigned char fused_mha_with_relPosBias_int8_64_32_kernel_sm75_cubin[];
extern unsigned char fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin[];
extern unsigned char fused_mha_with_relPosBias_int8_256_32_kernel_sm75_cubin[];
extern unsigned char fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin[];

extern unsigned char fused_multihead_attention_v2_int8_128_64_kernel_cubin[];
extern unsigned char fused_multihead_attention_v2_int8_192_64_kernel_cubin[];
extern unsigned char fused_multihead_attention_v2_int8_256_64_kernel_cubin[];
extern unsigned char fused_multihead_attention_v2_int8_384_64_kernel_cubin[];

extern unsigned char cubin_fmha_v2_int8_512_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_384_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_256_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_192_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_128_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_96_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_64_32_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_int8_512_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_384_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_256_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_192_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_128_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_96_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_64_64_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_int8_512_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_384_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_256_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_192_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_128_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_96_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_64_32_sm75_cu_cubin[];

extern unsigned char cubin_fmha_v2_int8_384_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_256_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_192_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_128_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_96_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_int8_64_64_sm75_cu_cubin[];

// NOTE: flash attention kernels for GPT (causal mask)
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin[];

extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin[];

extern unsigned int fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin_len;
extern unsigned int fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_32_sm75_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_64_sm75_cu_cubin_len;

extern unsigned int fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len;

extern unsigned int fused_mha_with_relPosBias_int8_64_32_kernel_sm75_cubin_len;
extern unsigned int fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin_len;
extern unsigned int fused_mha_with_relPosBias_int8_256_32_kernel_sm75_cubin_len;
extern unsigned int fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin_len;

extern unsigned int fused_multihead_attention_v2_int8_128_64_kernel_cubin_len;
extern unsigned int fused_multihead_attention_v2_int8_192_64_kernel_cubin_len;
extern unsigned int fused_multihead_attention_v2_int8_256_64_kernel_cubin_len;
extern unsigned int fused_multihead_attention_v2_int8_384_64_kernel_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin_len;

extern unsigned int cubin_fmha_v2_int8_512_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_384_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_256_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_192_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_128_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_96_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_64_32_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_int8_512_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_384_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_256_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_192_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_128_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_96_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_64_64_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_int8_512_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_384_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_256_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_192_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_128_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_96_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_64_32_sm75_cu_cubin_len;

extern unsigned int cubin_fmha_v2_int8_384_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_256_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_192_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_128_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_96_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_int8_64_64_sm75_cu_cubin_len;

static const struct FusedMultiHeadAttentionKernelMetaInfoV2 {
    Data_type            mDataType;
    unsigned int         mS;
    unsigned int         mD;
    unsigned int         mSM;
    const unsigned char* mCubin;
    unsigned int         mCubinSize;
    const char*          mFuncName;
    unsigned int         mSharedMemBytes;
    unsigned int         mThreadsPerCTA;
    unsigned int         mUnrollStep;
    bool                 mInterleaved;
    bool                 mWithRelativePositionBias;
    bool                 mFlashAttention = false;
    bool                 mCausalMask     = false;
} sMhaKernelMetaInfosV2[] = {
    // Xavier
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_128_64_kernel_cubin,
     fused_multihead_attention_v2_int8_128_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_128_64_kernel_sm72_interleaved",
     24576,
     128,
     0,
     true,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_128_64_kernel_cubin,
     fused_multihead_attention_v2_int8_128_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_128_64_kernel_sm72",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_192_64_kernel_cubin,
     fused_multihead_attention_v2_int8_192_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_192_64_kernel_sm72_interleaved",
     28672,
     128,
     0,
     true,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_192_64_kernel_cubin,
     fused_multihead_attention_v2_int8_192_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_192_64_kernel_sm72",
     45056,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_256_64_kernel_cubin,
     fused_multihead_attention_v2_int8_256_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_256_64_kernel_sm72_interleaved",
     36864,
     128,
     0,
     true,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_256_64_kernel_cubin,
     fused_multihead_attention_v2_int8_256_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_256_64_kernel_sm72",
     57344,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_384_64_kernel_cubin,
     fused_multihead_attention_v2_int8_384_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_384_64_kernel_sm72_interleaved",
     51200,
     128,
     0,
     true,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_72,
     fused_multihead_attention_v2_int8_384_64_kernel_cubin,
     fused_multihead_attention_v2_int8_384_64_kernel_cubin_len,
     "fused_multihead_attention_v2_int8_384_64_kernel_sm72",
     77824,
     128,
     0,
     false,
     false},
    // Volta
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_64_64_kernel_sm70",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_64_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_96_64_kernel_sm70",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_96_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_128_64_kernel_sm70",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_128_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_256_64_kernel_sm70",
     69632,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_256_64_kernel_sm70_noloop",
     69632,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_384_64_kernel_sm70",
     69632,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_384_64_kernel_sm70_noloop",
     69632,
     256,
     32,
     false,
     false},

    // Turing
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm75",
     12288,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm75_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm75",
     16384,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm75_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm75",
     18432,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm75_noloop",
     18432,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm75_kernel",
     28672,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm75_kernel_nl",
     28672,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm75_kernel",
     53248,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm75_kernel_nl",
     53248,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm75_kernel",
     14336,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm75_kernel_nl",
     14336,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm75_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm75_kernel_nl",
     28672,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm75_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm75_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm75_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm75_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_96_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_96_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm75_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_32_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm75_kernel",
     6144,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_32_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},

    {DATA_TYPE_INT8,
     64,
     32,
     kSM_75,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_int8_64_32_kernel_sm75",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_75,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm75_cubin,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm75_cubin_len,
     "fused_mha_with_relPosBias_int8_256_32_kernel_sm75",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_75,
     cubin_fmha_v2_int8_512_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm75_cu_cubin_len,
     "fmha_v2_int8_512_32_sm75_kernel",
     36864,
     256,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_75,
     cubin_fmha_v2_int8_512_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm75_cu_cubin_len,
     "fmha_v2_int8_512_32_sm75_kernel_nl",
     36864,
     256,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_int8_192_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm75_cu_cubin_len,
     "fmha_v2_int8_192_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_int8_192_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm75_cu_cubin_len,
     "fmha_v2_int8_192_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_int8_256_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm75_cu_cubin_len,
     "fmha_v2_int8_256_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_int8_256_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm75_cu_cubin_len,
     "fmha_v2_int8_256_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_int8_384_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm75_cu_cubin_len,
     "fmha_v2_int8_384_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_int8_384_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm75_cu_cubin_len,
     "fmha_v2_int8_384_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_int8_192_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm75_cu_cubin_len,
     "fmha_v2_int8_192_64_sm75_kernel",
     34816,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_int8_192_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm75_cu_cubin_len,
     "fmha_v2_int8_192_64_sm75_kernel_nl",
     34816,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_int8_256_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm75_cu_cubin_len,
     "fmha_v2_int8_256_64_sm75_kernel",
     34816,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_int8_256_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm75_cu_cubin_len,
     "fmha_v2_int8_256_64_sm75_kernel_nl",
     34816,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_int8_384_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm75_cu_cubin_len,
     "fmha_v2_int8_384_64_sm75_kernel",
     34816,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_int8_384_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm75_cu_cubin_len,
     "fmha_v2_int8_384_64_sm75_kernel_nl",
     34816,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_int8_128_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm75_cu_cubin_len,
     "fmha_v2_int8_128_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_int8_128_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm75_cu_cubin_len,
     "fmha_v2_int8_128_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_int8_128_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm75_cu_cubin_len,
     "fmha_v2_int8_128_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_int8_128_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm75_cu_cubin_len,
     "fmha_v2_int8_128_64_sm75_kernel_nl",
     34816,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     32,
     kSM_75,
     cubin_fmha_v2_int8_96_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_96_32_sm75_cu_cubin_len,
     "fmha_v2_int8_96_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     64,
     kSM_75,
     cubin_fmha_v2_int8_96_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_96_64_sm75_cu_cubin_len,
     "fmha_v2_int8_96_64_sm75_kernel",
     22528,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_int8_64_32_sm75_cu_cubin,
     cubin_fmha_v2_int8_64_32_sm75_cu_cubin_len,
     "fmha_v2_int8_64_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_int8_64_64_sm75_cu_cubin,
     cubin_fmha_v2_int8_64_64_sm75_cu_cubin_len,
     "fmha_v2_int8_64_64_sm75_kernel",
     36864,
     128,
     0,
     false,
     false},

#if CUDA_VERSION >= 11000
    // Ampere
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm80",
     12288,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm80_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm80",
     16384,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm80_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm80",
     18432,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm80_noloop",
     18432,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel",
     32768,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel_nl",
     32768,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel",
     65536,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel_nl",
     57344,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel_nl",
     16384,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel_nl",
     32768,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm80_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm80_kernel",
     49152,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm80_kernel",
     8192,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},

    {DATA_TYPE_INT8,
     64,
     32,
     kSM_80,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_int8_64_32_kernel_sm80",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_80,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_int8_256_32_kernel_sm80",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_80,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin_len,
     "fmha_v2_int8_512_32_sm80_kernel",
     40960,
     256,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_80,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin_len,
     "fmha_v2_int8_512_32_sm80_kernel_nl",
     40960,
     256,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     64,
     kSM_80,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin_len,
     "fmha_v2_int8_512_64_sm80_kernel",
     73728,
     256,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     64,
     kSM_80,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin_len,
     "fmha_v2_int8_512_64_sm80_kernel_nl",
     73728,
     256,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin_len,
     "fmha_v2_int8_192_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin_len,
     "fmha_v2_int8_192_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin_len,
     "fmha_v2_int8_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin_len,
     "fmha_v2_int8_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin_len,
     "fmha_v2_int8_384_32_sm80_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin_len,
     "fmha_v2_int8_384_32_sm80_kernel_nl",
     28672,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin_len,
     "fmha_v2_int8_192_64_sm80_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin_len,
     "fmha_v2_int8_192_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin_len,
     "fmha_v2_int8_256_64_sm80_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin_len,
     "fmha_v2_int8_256_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin_len,
     "fmha_v2_int8_384_64_sm80_kernel",
     53248,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin_len,
     "fmha_v2_int8_384_64_sm80_kernel_nl",
     53248,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin_len,
     "fmha_v2_int8_128_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin_len,
     "fmha_v2_int8_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin_len,
     "fmha_v2_int8_128_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin_len,
     "fmha_v2_int8_128_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     32,
     kSM_80,
     cubin_fmha_v2_int8_96_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_96_32_sm80_cu_cubin_len,
     "fmha_v2_int8_96_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     64,
     kSM_80,
     cubin_fmha_v2_int8_96_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_96_64_sm80_cu_cubin_len,
     "fmha_v2_int8_96_64_sm80_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_int8_64_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_64_32_sm80_cu_cubin_len,
     "fmha_v2_int8_64_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_int8_64_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_64_64_sm80_cu_cubin_len,
     "fmha_v2_int8_64_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},

    // GA10x
    // Note: For GA10X keep only kernels whose sharedMemBytes < 100KiB
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm86",
     12288,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_64_32_kernel_sm86_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm86",
     16384,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_128_32_kernel_sm86_noloop",
     10240,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm86",
     18432,
     128,
     0,
     false,
     true},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin,
     fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_cubin_len,
     "fused_mha_with_relPosBias_fp16_256_32_kernel_sm86_noloop",
     18432,
     128,
     32,
     false,
     true},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel",
     32768,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel_nl",
     32768,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel",
     65536,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel_nl",
     57344,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel_nl",
     16384,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel_nl",
     32768,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm80_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm80_kernel",
     49152,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm80_kernel",
     8192,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     32,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},

    {DATA_TYPE_INT8,
     64,
     32,
     kSM_86,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_int8_64_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_int8_64_32_kernel_sm80",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_86,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin,
     fused_mha_with_relPosBias_int8_256_32_kernel_sm80_cubin_len,
     "fused_mha_with_relPosBias_int8_256_32_kernel_sm80",
     10240,
     128,
     0,
     false,
     true},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_86,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin_len,
     "fmha_v2_int8_512_32_sm80_kernel",
     40960,
     256,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     32,
     kSM_86,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_32_sm80_cu_cubin_len,
     "fmha_v2_int8_512_32_sm80_kernel_nl",
     40960,
     256,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     64,
     kSM_86,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin_len,
     "fmha_v2_int8_512_64_sm80_kernel",
     73728,
     256,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     512,
     64,
     kSM_86,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_512_64_sm80_cu_cubin_len,
     "fmha_v2_int8_512_64_sm80_kernel_nl",
     73728,
     256,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin_len,
     "fmha_v2_int8_192_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_32_sm80_cu_cubin_len,
     "fmha_v2_int8_192_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin_len,
     "fmha_v2_int8_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_32_sm80_cu_cubin_len,
     "fmha_v2_int8_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin_len,
     "fmha_v2_int8_384_32_sm80_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_32_sm80_cu_cubin_len,
     "fmha_v2_int8_384_32_sm80_kernel_nl",
     28672,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin_len,
     "fmha_v2_int8_192_64_sm80_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_192_64_sm80_cu_cubin_len,
     "fmha_v2_int8_192_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin_len,
     "fmha_v2_int8_256_64_sm80_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_256_64_sm80_cu_cubin_len,
     "fmha_v2_int8_256_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin_len,
     "fmha_v2_int8_384_64_sm80_kernel",
     53248,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_384_64_sm80_cu_cubin_len,
     "fmha_v2_int8_384_64_sm80_kernel_nl",
     53248,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin_len,
     "fmha_v2_int8_128_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_32_sm80_cu_cubin_len,
     "fmha_v2_int8_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin_len,
     "fmha_v2_int8_128_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_128_64_sm80_cu_cubin_len,
     "fmha_v2_int8_128_64_sm80_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     32,
     kSM_86,
     cubin_fmha_v2_int8_96_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_96_32_sm80_cu_cubin_len,
     "fmha_v2_int8_96_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     96,
     64,
     kSM_86,
     cubin_fmha_v2_int8_96_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_96_64_sm80_cu_cubin_len,
     "fmha_v2_int8_96_64_sm80_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_int8_64_32_sm80_cu_cubin,
     cubin_fmha_v2_int8_64_32_sm80_cu_cubin_len,
     "fmha_v2_int8_64_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_INT8,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_int8_64_64_sm80_cu_cubin,
     cubin_fmha_v2_int8_64_64_sm80_cu_cubin_len,
     "fmha_v2_int8_64_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},
#endif
    // NOTE: flash attention
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm70_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm70_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm70_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm70_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm70_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm70_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm70_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm70_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm70_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm70_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm72_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm72_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm72_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm72_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm72_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm72_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm72_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm72_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm72_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm72_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm72_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm72_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm72_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm72_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm75_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm75_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm75_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm75_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm75_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm72_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm72_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},

    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm70_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm70_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm72_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_72,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm72_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm72_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm75_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm75_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     false},

    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm75_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm75_kernel",
     49152,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm75_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm72_kernel",
     49152,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm72_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm72_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm72_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm70_kernel",
     81920,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm70_kernel_nl",
     36864,
     256,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm70_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm70_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm70_kernel",
     36864,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm70_kernel_nl",
     36864,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm70_kernel",
     69632,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm70_kernel_nl",
     69632,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm72_kernel",
     36864,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm72_kernel_nl",
     36864,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm72_kernel",
     69632,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm72_kernel_nl",
     69632,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm70_kernel",
     20480,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm70_kernel_nl",
     18432,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm70_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm70_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm72_kernel",
     20480,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm72_kernel_nl",
     18432,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm72_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_72,
     cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm72_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm72_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},

#if CUDA_VERSION >= 11000
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm80_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm80_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm80_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm80_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm80_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm80_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm80_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm86_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm86_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm86_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm86_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm86_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm86_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm86_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm86_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm89_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm89_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm89_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm89_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm89_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm89_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm89_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm89_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},

    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm80_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm80_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm80_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm80_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm86_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_32_sm86_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm86_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     false},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_0_64_sm86_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     false},

    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm80_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm80_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm80_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm80_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm86_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm86_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm86_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm86_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm86_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm86_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm89_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm89_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm89_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm89_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm89_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm89_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true}
#endif
};

class FusedMultiHeadAttentionXMMAKernelV2:
    public TFusedMultiHeadAttentionXMMAKernel<FusedMultiHeadAttentionKernelMetaInfoV2,
                                              Fused_multihead_attention_params_v2> {
public:
    FusedMultiHeadAttentionXMMAKernelV2(const FusedMultiHeadAttentionKernelMetaInfoV2* pMetaStart,
                                        unsigned int                                   nMetaCount,
                                        Data_type                                      type,
                                        unsigned int                                   sm):
        TFusedMultiHeadAttentionXMMAKernel<FusedMultiHeadAttentionKernelMetaInfoV2,
                                           Fused_multihead_attention_params_v2>(pMetaStart, nMetaCount, type, sm)
    {
    }

    inline uint64_t hashID(unsigned int s,
                           unsigned int d,
                           bool         interleaved,
                           bool         unroll,
                           bool         with_relative_position_bias,
                           bool         flash_attention = false,
                           bool         causal_mask     = false) const
    {
        s = flash_attention ? 0 : s;
        return (uint64_t)s << 32 | d << 8 | (with_relative_position_bias ? 16ull : 0ull)
               | (flash_attention ? 8ull : 0ull) | (causal_mask ? 4ull : 0ull) | (interleaved ? 2ull : 0ull)
               | (unroll ? 1ull : 0ull);
    }

    virtual uint64_t hashID(const KernelMeta& kernelMeta) const
    {
        assert(kernelMeta.mD == 64 || kernelMeta.mD == 32 || kernelMeta.mD == 40 || kernelMeta.mD == 80 || kernelMeta.mD == 128
               || kernelMeta.mD == 144 || kernelMeta.mD == 160 || kernelMeta.mD == 256);
        return hashID(kernelMeta.mS,
                      kernelMeta.mD,
                      kernelMeta.mInterleaved,
                      kernelMeta.mUnrollStep,
                      kernelMeta.mWithRelativePositionBias,
                      kernelMeta.mFlashAttention,
                      kernelMeta.mCausalMask);
    }

    virtual void run(Fused_multihead_attention_params_v2& params,
                     cudaStream_t                         ss,
                     bool                                 flash_attention = false,
                     bool                                 causal_mask     = false) const
    {
        // assert(params.d == 64 || params.d == 32);
        if (params.interleaved) {
            assert(mDataType == fastertransformer::DATA_TYPE_INT8);
        }

        bool forceUnroll = params.force_unroll;
        if (!forceUnroll && !params.ignore_b1opt && mSM >= kSM_75) {
            const struct {
                unsigned int mSM;
                Data_type    mDataType;
                int          mS;
                int          mD;
                int          mMaxBatch;
            } unrollList[] = {
                {kSM_75, fastertransformer::DATA_TYPE_FP16, 128, 32, 1},
                {kSM_75, fastertransformer::DATA_TYPE_FP16, 128, 40, 1},
                {kSM_75, fastertransformer::DATA_TYPE_FP16, 128, 64, 1},
                {kSM_75, fastertransformer::DATA_TYPE_FP16, 256, 64, 1},
                {kSM_75, fastertransformer::DATA_TYPE_FP16, 384, 64, 1},
                {kSM_75, fastertransformer::DATA_TYPE_INT8, 128, 64, 1},
                {kSM_75, fastertransformer::DATA_TYPE_INT8, 192, 64, 2},
                {kSM_75, fastertransformer::DATA_TYPE_INT8, 256, 64, 1},
                {kSM_75, fastertransformer::DATA_TYPE_INT8, 384, 64, 1},
#if CUDA_VERSION >= 11000
                {kSM_80, fastertransformer::DATA_TYPE_INT8, 192, 64, 16},
                {kSM_80, fastertransformer::DATA_TYPE_INT8, 256, 64, 8},
                {kSM_80, fastertransformer::DATA_TYPE_INT8, 384, 64, 8},

                {kSM_86, fastertransformer::DATA_TYPE_FP16, 128, 32, 4},
                {kSM_86, fastertransformer::DATA_TYPE_FP16, 128, 40, 4},
                {kSM_86, fastertransformer::DATA_TYPE_FP16, 128, 64, 4},
                {kSM_86, fastertransformer::DATA_TYPE_FP16, 256, 64, 4},
                {kSM_86, fastertransformer::DATA_TYPE_INT8, 128, 64, 4},
                {kSM_86, fastertransformer::DATA_TYPE_INT8, 192, 64, 16},
                {kSM_86, fastertransformer::DATA_TYPE_INT8, 256, 64, 8},
                {kSM_86, fastertransformer::DATA_TYPE_INT8, 384, 64, 8},
#endif
            };
            for (unsigned int i = 0u; i < sizeof(unrollList) / sizeof(unrollList[0]); ++i) {
                if (mSM == unrollList[i].mSM && mDataType == unrollList[i].mDataType && params.s == unrollList[i].mS
                    && params.d == unrollList[i].mD && params.b <= unrollList[i].mMaxBatch) {
                    forceUnroll = true;
                    break;
                }
            }
        }

        bool       withRelativePositionBias = params.window_num == 0 ? false : true;
        const auto findIter                 = mFunctions.find(hashID(params.s,
                                                     params.d,
                                                     params.interleaved,
                                                     forceUnroll,
                                                     withRelativePositionBias,
                                                     flash_attention,
                                                     causal_mask));
        assert(findIter != mFunctions.end());

        const auto&      kernelMeta     = mKernelMeta[findIter->second.mMetaInfoIndex];
        const CUfunction func           = findIter->second.mDeviceFunction;
        void*            kernelParams[] = {&params, nullptr};
        if (!forceUnroll) {
            cuErrCheck(mDriver.cuLaunchKernel(func,
                                              params.h,
                                              params.b,
                                              1,
                                              kernelMeta.mThreadsPerCTA,
                                              1,
                                              1,
                                              kernelMeta.mSharedMemBytes,
                                              ss,
                                              kernelParams,
                                              nullptr),
                       mDriver);
        }
        else {
            int unroll = kernelMeta.mS / kernelMeta.mUnrollStep;
            assert(kernelMeta.mS == kernelMeta.mUnrollStep * unroll);
            // flash attention supports any sequence length, so we runtime s here
            if (flash_attention) {
                unroll = (params.s + kernelMeta.mUnrollStep - 1) / kernelMeta.mUnrollStep;
            }
            cuErrCheck(mDriver.cuLaunchKernel(func,
                                              params.h,
                                              params.b,
                                              unroll,
                                              kernelMeta.mThreadsPerCTA,
                                              1,
                                              1,
                                              kernelMeta.mSharedMemBytes,
                                              ss,
                                              kernelParams,
                                              nullptr),
                       mDriver);
        }
    }
};

using FusedMHAKernelFactoryV2 = TFusedMHAKernelFactory<FusedMultiHeadAttentionXMMAKernelV2>;

inline const FusedMultiHeadAttentionXMMAKernelV2* getXMMAKernelsV2(Data_type type, unsigned int sm)
{
    return FusedMHAKernelFactoryV2::Get().getXMMAKernels(
        sMhaKernelMetaInfosV2, sizeof(sMhaKernelMetaInfosV2) / sizeof(sMhaKernelMetaInfosV2[0]), type, sm);
}

}  // namespace fastertransformer
