/**
 * 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.
 */

#ifndef VERIFY_SINGLE_HEADER
#include "kernel_operator.h"
#else
#include "kernel_common.h"
#endif
#include "kernel_tensor.h" // todo

// __aicore__ static inline void SetVectorMask(const uint64_t maskHigh, const uint64_t maskLow)
extern "C" __global__ __aicore__ void KernelTestSetVectorMask1()
{
    uint64_t maskHigh = 0xFFFFFFFFFFFFFFFF;
    uint64_t maskLow = 0xFFFFFFFFFFFFFFFF;
    AscendC::SetVectorMask<uint32_t>(maskHigh, maskLow);
}

// __aicore__ static inline void SetVectorMask(int32_t len)
extern "C" __global__ __aicore__ void KernelTestSetVectorMask2()
{
    int32_t len = 0;
    AscendC::SetVectorMask<uint32_t>(len);
}

// __aicore__ inline void ResetMask()
extern "C" __global__ __aicore__ void KernelTestResetMask1() { AscendC::ResetMask(); }

// __aicore__ inline void SetMaskCount()
extern "C" __global__ __aicore__ void KernelTestSetMaskCount1() { AscendC::SetMaskCount(); }

// __aicore__ inline void SetMaskNorm()
extern "C" __global__ __aicore__ void KernelTestSetMaskNorm1() { AscendC::SetMaskNorm(); }

// __aicore__ inline void SetHcclContext(__gm__ uint8_t* context)
extern "C" __global__ __aicore__ void KernelTestSetHcclContext1()
{
    __gm__ uint8_t* context;
    AscendC::SetHcclContext<1>(context);
}

#if defined(__NPU_ARCH__) && ((__NPU_ARCH__ == 2201) || (__NPU_ARCH__ == 2002) || (__NPU_ARCH__ == 3002) || \
                              (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102))
// __aicore__ inline void SetAippFunctions(const GlobalTensor<T>& src0, AippInputFormat format, AippParams<U> config)
extern "C" __global__ __aicore__ void KernelTestSetAippFunctions1()
{
    AscendC::GlobalTensor<uint8_t> src0;
    AscendC::AippInputFormat format;
    AscendC::AippParams<uint8_t> config;
    AscendC::SetAippFunctions(src0, format, config);
}

// __aicore__ inline void SetAippFunctions(const GlobalTensor<T>& src0, const GlobalTensor<T>& src1, AippInputFormat
// format, AippParams<U> config)
extern "C" __global__ __aicore__ void KernelTestSetAippFunctions2()
{
    AscendC::GlobalTensor<uint8_t> src0;
    AscendC::GlobalTensor<uint8_t> src1;
    AscendC::AippInputFormat format;
    AscendC::AippParams<uint8_t> config;
    AscendC::SetAippFunctions(src0, src1, format, config);
}
#endif