/**  * Copyright (c) 2026 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 compressor.cpp * \brief */ #if (__CCE_AICORE__ == 220) #include "arch32/compressor_kernel.h" #include "arch32/compressor_kernel_perf.h" #else #include "arch35/compressor_kernel.h" #include "arch35/compressor_kernel_full_load.h" #endif using namespace Compressor; #define INVOKE_COMPRESSOR_GENERAL_OP_IMPL(templateClass, ...) \ do { \ templateClass> op(&pipe, tilingData); \ op.Init(x, wKv, wGate, stateCache, ape, normWeight, ropeSin, ropeCos, stateBlockTable, \ cuSeqlens, seqUsed, startPos, cmpKvOut, workspace); \ op.Process(); \ } while (0) #if (__CCE_AICORE__ == 220) template #else template #endif __global__ __aicore__ void compressor( __gm__ uint8_t *x, __gm__ uint8_t *wKv, __gm__ uint8_t *wGate, __gm__ uint8_t *stateCache, __gm__ uint8_t *ape, __gm__ uint8_t *normWeight, __gm__ uint8_t *ropeSin, __gm__ uint8_t *ropeCos, __gm__ uint8_t *stateBlockTable, __gm__ uint8_t *cuSeqlens, __gm__ uint8_t *seqUsed, __gm__ uint8_t *startPos, __gm__ uint8_t *cmpKvOut, __gm__ uint8_t *stateCacheOut, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { REGISTER_TILING_DEFAULT(optiling::CompressorTilingData); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2); GET_TILING_DATA_WITH_STRUCT(optiling::CompressorTilingData, tilingDataIn, tiling); if constexpr (static_cast(TemplateId) == TEMPLATE_ID::EMPTY_X) { return; } const optiling::CompressorTilingData *__restrict tilingData = &tilingDataIn; TPipe pipe; constexpr auto xLayout = static_cast(XLayout); constexpr auto xDtype = static_cast(XDType); #if (__CCE_AICORE__ == 220) constexpr auto ropeDtype = static_cast(RopeDType); #endif constexpr auto coff = static_cast(Coff); constexpr auto rotaryMode = static_cast(RotaryMode); #if (__CCE_AICORE__ != 220) constexpr auto cacheMode = static_cast(CacheMode); #endif #if (__CCE_AICORE__ == 220) if constexpr (static_cast(TemplateId) == TEMPLATE_ID::PERF) { INVOKE_COMPRESSOR_GENERAL_OP_IMPL(CompressorKernelPerf, xLayout, xDtype, ropeDtype, coff, rotaryMode); } else { INVOKE_COMPRESSOR_GENERAL_OP_IMPL(CompressorKernel, xLayout, xDtype, ropeDtype, coff, rotaryMode); } #else if constexpr (static_cast(TemplateId) == TEMPLATE_ID::FULL_LOAD) { INVOKE_COMPRESSOR_GENERAL_OP_IMPL(CompressorKernelFullLoad, xLayout, xDtype, coff, rotaryMode, cacheMode); } else { INVOKE_COMPRESSOR_GENERAL_OP_IMPL(CompressorKernel, xLayout, xDtype, coff, rotaryMode, cacheMode); } #endif }