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

/* Generated By CANNBot */

/*!
 * \file threshold_simt.h
 * \brief SIMT kernel implementation for threshold
 */

#ifndef THRESHOLD_SIMT_H
#define THRESHOLD_SIMT_H

#include "kernel_operator.h"
#include "kernel_tiling/kernel_tiling.h"
#include "simt_api/common_functions.h"
#include "simt_api/asc_simt.h"
#include "threshold_tiling_data.h"
#include "threshold_tiling_key.h"

namespace NsThreshold {

using namespace AscendC;

constexpr uint32_t THREAD_NUM = 512;

template <typename T>
__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void ThresholdSimtKernel(int64_t totalElements,
                                                                                     float threshold, __gm__ T* input,
                                                                                     __gm__ T* output)
{
    // Output constants in target dtype (0.0 and 1.0 are exactly representable in all FP formats)
    const T one = static_cast<T>(1.0f);
    const T zero = static_cast<T>(0.0f);

    // Grid-Stride loop: each thread traverses all elements
    for (uint64_t idx = static_cast<uint64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         idx < static_cast<uint64_t>(totalElements); idx += static_cast<uint64_t>(blockDim.x) * gridDim.x) {
        // Upcast to float32 for comparison, matching golden semantics
        // (golden converts input to float32 and compares with float32 threshold)
        float valF32 = static_cast<float>(input[idx]);
        output[idx] = (valF32 > threshold) ? one : zero;
    }
}

template <typename T>
__aicore__ inline void Process(GM_ADDR input, GM_ADDR output, const ThresholdTilingData* tilingData)
{
    int64_t totalElements = tilingData->totalElements;
    float threshold = tilingData->threshold;

    __gm__ T* inputGm = (__gm__ T*)input;
    __gm__ T* outputGm = (__gm__ T*)output;

    asc_vf_call<ThresholdSimtKernel<T>>(dim3(THREAD_NUM), totalElements, threshold, inputGm, outputGm);
}

} // namespace NsThreshold

#endif