Star0
0
代码
代码
Issues
Pull Requests
Actions
项目讨论
Wiki
分析
Star0
0
cann-robotcann-robot修改DequantSwigluQuant算子，使用饱和模式
3c60192e创建于 5月7日历史提交
/**
 * 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
}