* This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
* and is contributed to the CANN Open Software.
*
* Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
* All Rights Reserved.
*
* Authors (accounts):
* - Zhou Jiamin <@zhou-jiamin-666>
* - Su Tonghua <@sutonghua>
*
* This program is free software: you can redistribute it and/or modify it.
* Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
* You may not use this file except in compliance with the License.
* See the LICENSE file at the root of the repository for the full text of the License.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
*/
* \file lin_space_d.cpp
* \brief
*/
#include "lin_space_d.h"
using namespace NsLinSpaceD;
enum class LinSpaceDTilingKey : uint32_t
{
TILING_KEY_TPL_BF16 = 27,
TILING_KEY_TPL_UINT8 = 4,
TILING_KEY_TPL_FP16 = 1,
TILING_KEY_TPL_FP32 = 0,
TILING_KEY_TPL_INT8 = 2,
TILING_KEY_TPL_INT16 = 6,
TILING_KEY_TPL_INT32 = 3,
};
using BF16 = bfloat16_t;
using UINT8 = uint8_t;
using FP16 = half;
using FP32 = float;
using INT8 = int8_t;
using INT16 = int16_t;
using INT32 = int32_t;
template <uint32_t TYPE_START, uint32_t TYPE_END>
__global__ __aicore__ void lin_space_d(GM_ADDR start, GM_ADDR end, GM_ADDR num, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
if (workspace == nullptr) {
return;
}
REGISTER_TILING_DEFAULT(LinSpaceDTilingData);
GET_TILING_DATA_WITH_STRUCT(LinSpaceDTilingData, tilingData, tiling);
if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<BF16, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<BF16, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<BF16, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<BF16, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<BF16, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<BF16, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<BF16, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<UINT8, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<UINT8, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<UINT8, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<UINT8, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<UINT8, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<UINT8, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<UINT8, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<FP16, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<FP16, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<FP16, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<FP16, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<FP16, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<FP16, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<FP16, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<FP32, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<FP32, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<FP32, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<FP32, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<FP32, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<FP32, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<FP32, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<INT8, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<INT8, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<INT8, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<INT8, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<INT8, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<INT8, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<INT8, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<INT16, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<INT16, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<INT16, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<INT16, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<INT16, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<INT16, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<INT16, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
else if constexpr (TYPE_START == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_BF16)) {
LinSpaceD<INT32, BF16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_UINT8)) {
LinSpaceD<INT32, UINT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP16)) {
LinSpaceD<INT32, FP16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_FP32)) {
LinSpaceD<INT32, FP32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT8)) {
LinSpaceD<INT32, INT8> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT16)) {
LinSpaceD<INT32, INT16> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
else if constexpr (TYPE_END == static_cast<uint32_t>(LinSpaceDTilingKey::TILING_KEY_TPL_INT32)) {
LinSpaceD<INT32, INT32> op;
op.Init(start, end, z, &tilingData);
op.Process();
}
}
}