Star813
1612
代码
代码
Issues158
Pull Requests234
Actions
项目讨论
Wiki
分析
Star813
1612
cann-robotcann-robotformat cpp
5db91519创建于 22 天前历史提交
/**
 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
 * CANN Open Software License Agreement Version 2.0 (the "License").
 * Please refer to the License for details. You may not use this file except in compliance with the License.
 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
 * See LICENSE in the root of the software repository for the full text of the License.
 */

/*!
 * \file dequant_swiglu_quant_apt.cpp
 * \brief
 */
#if ASC_DEVKIT_MAJOR >= 9
#include "basic_api/kernel_vec_intf.h"
#else
#include "kernel_operator.h"
#endif
#include "arch35/dequant_swiglu_quant.h"
#include "arch35/dequant_swiglu_quant_nlast.h"
#include "arch35/dequant_swiglu_quant_static.h"
#include "arch35/dequant_swiglu_quant_dynamic_not_full.h"
#include "arch35/dequant_swiglu_quant_static_not_full.h"

#define ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE 111
#define ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 110
#define ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE 101
#define ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 100
#define ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE 11
#define ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 10
#define ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE 1
#define ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 0
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE 1111    // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE 2111    // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE 3111    // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE 4111    // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 1110   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 2110   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 3110   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 4110   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE 1101   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE 2101   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE 3101   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE 4101   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 1100  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 2100  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 3100  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 4100  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE 1011   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE 2011   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE 3011   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE 4011   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 1010  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 2010  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 3010  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 4010  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE 1001  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE 2001  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE 3001  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE 4001  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 1000 // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 2000 // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 3000 // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 4000 // bias = bfloat16
#define ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 100110
#define ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 100100
#define ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 100010
#define ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 100000
#define ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 101110   // bias = int32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 102110   // bias = float32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 103110   // bias = float16
#define ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE 104110   // bias = bfloat16
#define ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 101100  // bias = int32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 102100  // bias = float32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 103100  // bias = float16
#define ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE 104100  // bias = bfloat16
#define ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 101010  // bias = int32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 102010  // bias = float32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 103010  // bias = float16
#define ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE 104010  // bias = bfloat16
#define ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 101000 // bias = int32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 102000 // bias = float32
#define ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 103000 // bias = float16
#define ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE 104000 // bias = bfloat16
// static quant not full load
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1000100
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE 1000110
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1001100
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE 1001110
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1010100
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE 1010110
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1011100
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE 1011110
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1000101
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE 1000111
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1001101
#define QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE 1001111
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1010101
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE 1010111
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1011101
#define QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE 1011111

#define ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 10111
#define ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 10110
#define ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 10101
#define ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 10100
#define ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 10011
#define ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 10010
#define ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 10001
#define ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 10000
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 11111    // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 12111    // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 13111    // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 14111    // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 11110   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 12110   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 13110   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 14110   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 11101   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 12101   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 13101   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 14101   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 11100  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 12100  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 13100  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 14100  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 11011   // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 12011   // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 13011   // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC 14011   // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 11010  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 12010  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 13010  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC 14010  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 11001  // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 12001  // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 13001  // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC 14001  // bias = bfloat16
#define ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 11000 // bias = int32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 12000 // bias = float32
#define ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 13000 // bias = float16
#define ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC 14000 // bias = bfloat16
// ##Dynamic Not Full Load
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1100000
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1100001
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1100100
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1100101
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1101000
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1101001
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1101100
#define QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1101101
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1110000
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1110001
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1110100
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1110101
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1111000
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1111001
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1111100
#define QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1111101
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1120000
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1120001
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1120100
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1120101
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1121000
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1121001
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1121100
#define QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1121101
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1130000
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1130001
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1130100
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1130101
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1131000
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1131001
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1131100
#define QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1131101
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1140000
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1140001
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1140100
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1140101
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE 1141000
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE 1141001
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE 1141100
#define QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE 1141101

#ifndef GLOBAL_OVERFLOW_MODE_CTRL
#define GLOBAL_OVERFLOW_MODE_CTRL 60
#endif

using namespace AscendC;

extern "C" __global__ __aicore__ void dequant_swiglu_quant(GM_ADDR x, GM_ADDR weightScale, GM_ADDR activationScale,
                                                           GM_ADDR bias, GM_ADDR quantScale, GM_ADDR quantOffset,
                                                           GM_ADDR groupIndex, GM_ADDR y, GM_ADDR scale,
                                                           GM_ADDR workspace, GM_ADDR tiling)
{
    if (g_coreType != AIV) {
        return;
    }

    GM_ADDR globalWorkGm = GetUserWorkspace(workspace);
    TPipe pipe;
    TILING_KEY_IS(ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    TILING_KEY_IS(ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE);
    // static quant not full load
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE);

    TILING_KEY_IS(ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    TILING_KEY_IS(ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC);
    // ##Dynamic Not Full Load
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE);
    TILING_KEY_IS(QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE);

#if (__NPU_ARCH__ == 3510)
    int64_t oriOverflowMode = AscendC::GetCtrlSpr<GLOBAL_OVERFLOW_MODE_CTRL, GLOBAL_OVERFLOW_MODE_CTRL>();
#endif

#if TILING_KEY_VAR == ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale有值, group_index有值, 000111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale有值, group_index无值, 000110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale无值, group_index有值, 000101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale无值, group_index无值, 000100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale有值, group_index有值, 000011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale有值, group_index无值, 000010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale无值, group_index有值, 000001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale无值, group_index无值, 000000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale有值, group_index有值, 001111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale有值, group_index有值, 002111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale有值, group_index有值, 003111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale有值, group_index有值, 004111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale有值, group_index无值, 001110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale有值, group_index无值, 002110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale有值, group_index无值, 003110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale有值, group_index无值, 004110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale无值, group_index有值, 001101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale无值, group_index有值, 002101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale无值, group_index有值, 003101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale无值, group_index有值, 004101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale无值, group_index无值, 001100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale无值, group_index无值, 002100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale无值, group_index无值, 003100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale无值, group_index无值, 004100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<float, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale有值, group_index有值, 001011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale有值, group_index有值, 002011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale有值, group_index有值, 003011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale有值, group_index有值, 004011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale有值, group_index无值, 001010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale有值, group_index无值, 002010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale有值, group_index无值, 003010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale有值, group_index无值, 004010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale无值, group_index有值, 001001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale无值, group_index有值, 002001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale无值, group_index有值, 003001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale无值, group_index有值, 004001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale无值, group_index无值, 001000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale无值, group_index无值, 002000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale无值, group_index无值, 003000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale无值, group_index无值, 004000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBase<bool, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale有值, group_index无值, 101110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale有值, group_index无值, 102110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale有值, group_index无值, 103110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale有值, group_index无值, 104110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale无值, group_index无值, 101100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale无值, group_index无值, 102100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=flaot16, activate_scale有值, quant_scale无值, group_index无值, 103100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale无值, group_index无值, 104100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale有值, group_index无值, 101010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale有值, group_index无值, 102010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale有值, group_index无值, 103010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale有值, group_index无值, 104010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale无值, group_index无值, 101000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=flaot32, activate_scale无值, quant_scale无值, group_index无值, 102000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale无值, group_index无值, 103000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale无值, group_index无值, 104000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale无值, group_index无值, 100000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale有值, group_index无值, 100010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<bool, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale无值, group_index无值, 100100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_TRUE_X_FALSE_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE
    // activate_dim!=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale有值, group_index无值, 100110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35NlastTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35NlastTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantNlast<float, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale有值, group_index有值, 000111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale有值, group_index无值, 000110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale无值, group_index有值, 000101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale有值, quant_scale无值, group_index无值, 000100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale有值, group_index有值, 000011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale有值, group_index无值, 000010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale无值, group_index有值, 000001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, int64_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias无值, activate_scale无值, quant_scale无值, group_index无值, 000000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale有值, group_index有值, 001111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale有值, group_index有值, 002111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale有值, group_index有值, 003111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale有值, group_index有值, 004111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale有值, group_index无值, 001110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale有值, group_index无值, 002110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale有值, group_index无值, 003110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale有值, group_index无值, 004110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale无值, group_index有值, 001101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale无值, group_index有值, 002101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale无值, group_index有值, 003101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale无值, group_index有值, 004101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale有值, quant_scale无值, group_index无值, 001100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale有值, quant_scale无值, group_index无值, 002100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale有值, quant_scale无值, group_index无值, 003100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_TRUE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale有值, quant_scale无值, group_index无值, 004100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<float, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale有值, group_index有值, 001011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale有值, group_index有值, 002011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale有值, group_index有值, 003011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale有值, group_index有值, 004011
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale有值, group_index无值, 001010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale有值, group_index无值, 002010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale有值, group_index无值, 003010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_TRUE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale有值, group_index无值, 004010
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale无值, group_index有值, 001001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, int64_t, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale无值, group_index有值, 002001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, int64_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale无值, group_index有值, 003001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, int64_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_TRUE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale无值, group_index有值, 004001
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, int64_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_TRUE_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=int32, activate_scale无值, quant_scale无值, group_index无值, 001000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP32_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float32, activate_scale无值, quant_scale无值, group_index无值, 002000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_FP16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=float16, activate_scale无值, quant_scale无值, group_index无值, 003000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == ACTDIM_FALSE_X_FALSE_BIAS_BF16_ACT_FALSE_QUANT_FALSE_GROUPINDEX_FALSE_STATIC
    // activate_dim=-1, x=int32/fp16/bf6, bias=bfloat16, activate_scale无值, quant_scale无值, group_index无值, 004000
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantBaseStatic<bool, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, int32_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, int32_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, int32_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, bool, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, int32_t, bool, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, int32_t, int32_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, int32_t, int32_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, int32_t, int32_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, bool, int32_t, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_INT32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, int32_t, int32_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, int32_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, int32_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, int32_t, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, bool, float, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP32_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, int32_t, float, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, int32_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, int32_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, int32_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, bool, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_FP16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, int32_t, half, DTYPE_X, DTYPE_Y> op(&pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, bool, int32_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<bool, float, int32_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_FALSE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, bool, int32_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, bool, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_DYNAMIC_BIAS_BF16_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantDynamicNotFull<float, float, int32_t, bfloat16_t, DTYPE_X, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, globalWorkGm, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    // 静态量化, bias无值, activate_scale无值, quant_scale有值, quant_offset无值, 无group, 1000100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, bool, float, bool, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, nullptr, quantScale, nullptr, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE
    // 静态量化, bias无值, activate_scale无值, quant_scale有值, quant_offset有值, 无group, 1000110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, bool, float, float, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, nullptr, quantScale, quantOffset, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    // 静态量化, bias无值, activate_scale有值, quant_scale有值, quant_offset无值, 无group, 1001100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, bool, float, bool, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE
    // 静态量化, bias无值, activate_scale有值, quant_scale有值, quant_offset有值, 无group, 1001110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, bool, float, float, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    // 静态量化, bias有值, activate_scale无值, quant_scale有值, quant_offset无值, 无group, 1010100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, float, float, bool, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, bias, quantScale, nullptr, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE
    // 静态量化, bias有值, activate_scale无值, quant_scale有值, quant_offset有值, 无group, 1010110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, float, float, float, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, bias, quantScale, quantOffset, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_FALSE
    // 静态量化, bias有值, activate_scale有值, quant_scale有值, quant_offset无值, 无group, 1011100
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, float, float, bool, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_FALSE
    // 静态量化, bias有值, activate_scale有值, quant_scale有值, quant_offset有值, 无group, 1011110
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, float, float, float, bool, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, nullptr, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    // 静态量化, bias无值, activate_scale无值, quant_scale有值, quant_offset无值, 有group, 1000101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, bool, float, bool, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE
    // 静态量化, bias无值, activate_scale无值, quant_scale有值, quant_offset有值, 有group, 1000111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, bool, float, float, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    // 静态量化, bias无值, activate_scale有值, quant_scale有值, quant_offset无值, 有group, 1001101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, bool, float, bool, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_FALSE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE
    // 静态量化, bias无值, activate_scale有值, quant_scale有值, quant_offset有值, 有group, 1001111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, bool, float, float, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, nullptr, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    // 静态量化, bias有值, activate_scale无值, quant_scale有值, quant_offset无值, 有group, 1010101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, float, float, bool, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_FALSE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE
    // 静态量化, bias有值, activate_scale无值, quant_scale有值, quant_offset有值, 有group, 1010111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, bool, float, float, float, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, nullptr, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_FALSE_GROUP_TRUE
    // 静态量化, bias有值, activate_scale有值, quant_scale有值, quant_offset无值, 有group, 1011101
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, float, float, bool, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, nullptr, groupIndex, y, scale, tilingData);
    op.Process();
#elif TILING_KEY_VAR == QUANTMODE_STATIC_BIAS_TRUE_ACT_TRUE_QUANT_TRUE_OFFSET_TRUE_GROUP_TRUE
    // 静态量化, bias有值, activate_scale有值, quant_scale有值, quant_offset有值, 有group, 1011111
    GET_TILING_DATA_WITH_STRUCT(DequantSwigluQuantV35BaseTilingData, tilingDataIn, tiling);
    const DequantSwigluQuantV35BaseTilingData* __restrict tilingData = &tilingDataIn;
    DequantSwigluQuantV35Ops::DequantSwigluQuantStaticNotFull<DTYPE_X, float, float, float, float, int32_t, DTYPE_Y> op(
        &pipe);
    op.Init(x, weightScale, activationScale, bias, quantScale, quantOffset, groupIndex, y, scale, tilingData);
    op.Process();
#endif
#if (__NPU_ARCH__ == 3510)
    AscendC::SetCtrlSpr<GLOBAL_OVERFLOW_MODE_CTRL, GLOBAL_OVERFLOW_MODE_CTRL>(oriOverflowMode);
#endif
}