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

/**
 * NOTE: Portions of this code were AI-generated and have been
 * technically reviewed for functional accuracy and security
 * Generated By CANNBot
 */

/*!
 * \file apply_proximal_adagrad.cpp
 * \brief ApplyProximalAdagrad kernel entry (arch35 / Ascend950).
 *
 * Signature follows registry-invoke convention:
 *   all inputs -> all outputs -> workspace -> tiling
 *
 * Template parameters are driven by ASCENDC_TPL_SEL_PARAM on the host:
 *   - D_T_VAR  : input dtype (C_DT_FLOAT / C_DT_FLOAT16 / C_DT_BF16).
 *   - PAD_TAIL : whether the per-core tail is non-32B-aligned (0/1).
 *                In the current build the kernel always uses DataCopyPad on
 *                the trailing tile, so the parameter is reserved for future
 *                aligned-only fast-path optimisations.
 *   - HAS_L1   : compile-time hint that l1 may be non-zero. When 0, the
 *                kernel statically drops the sign + soft-threshold branch.
 *
 * fp16 / bf16 inputs are cast to fp32 inside the kernel template
 * (see arch35/apply_proximal_adagrad.h Compute()) and cast back on output.
 */

#include "./arch35/apply_proximal_adagrad.h"

#ifdef __CCE_KT_TEST__
// CPU UT 入口：保留单一非模板 extern "C" 符号，dtype 通过编译期宏
// `DTYPE_VAR` 注入（见 tests/ut/op_kernel/CMakeLists.txt 的 `AddOpTestCase`）。
// HAS_L1 固化为 true，使 UT 走通用分支；PAD_TAIL 固化为 false，因 UT 用例
// 选用 32B 对齐的元素数。
//
// 算子对外只暴露 1 个输出 var_out, accum 通过 input ref 原地更新, 因此
// kernel 入口签名为 (var, accum, lr, l1, l2, grad, var_out, workspace, tiling),
// 内部把 accum 形参既当作输入又当作 accumOut 转给 Init。
extern "C" __global__ __aicore__ void apply_proximal_adagrad(GM_ADDR var, GM_ADDR accum, GM_ADDR lr, GM_ADDR l1,
                                                             GM_ADDR l2, GM_ADDR grad, GM_ADDR var_out,
                                                             GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA_WITH_STRUCT(ApplyProximalAdagradTilingData, tilingData, tiling);
    NsApplyProximalAdagrad::ApplyProximalAdagrad<DTYPE_VAR, false, true> op;
    op.Init(var, accum, lr, l1, l2, grad, var_out, accum, &tilingData);
    op.Process();
}
#else
template <typename D_T_VAR, uint32_t PAD_TAIL, uint32_t HAS_L1>
__global__ __aicore__ void apply_proximal_adagrad(GM_ADDR var, GM_ADDR accum, GM_ADDR lr, GM_ADDR l1, GM_ADDR l2,
                                                  GM_ADDR grad, GM_ADDR var_out, GM_ADDR workspace, GM_ADDR tiling)
{
    REGISTER_TILING_DEFAULT(ApplyProximalAdagradTilingData);
    GET_TILING_DATA_WITH_STRUCT(ApplyProximalAdagradTilingData, tilingData, tiling);
    NsApplyProximalAdagrad::ApplyProximalAdagrad<D_T_VAR, (PAD_TAIL != 0U), (HAS_L1 != 0U)> op;
    // accum 既作为输入也作为 inplace 写回端口, 转给 Init 的 accumOut 形参。
    op.Init(var, accum, lr, l1, l2, grad, var_out, accum, &tilingData);
    op.Process();
}
#endif