/**
* Copyright (c) 2025 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 matmul_leakyrelu.asc
* \brief
*/
#include "data_utils.h"
#include "kernel_tiling/kernel_tiling.h"
#include "tiling/platform/platform_ascendc.h"
#include "tiling/tiling_api.h"
#include "acl/acl.h"
#include "kernel_operator.h"
#include "lib/matmul_intf.h"
__aicore__ inline uint32_t Ceiling(uint32_t a, uint32_t b)
{
return (a + b - 1) / b;
}
template <typename aType, typename bType, typename cType, typename biasType> class MatmulLeakyKernel {
public:
__aicore__ inline MatmulLeakyKernel(){};
__aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace,
const TCubeTiling &tiling, AscendC::TPipe *pipe);
__aicore__ inline void Process(AscendC::TPipe *pipe);
__aicore__ inline void MatmulCompute();
__aicore__ inline void LeakyReluCompute();
__aicore__ inline void CopyOut(uint32_t count);
__aicore__ inline void CalcOffset(int32_t blockIdx, const TCubeTiling &tiling, int32_t &offsetA, int32_t &offsetB,
int32_t &offsetC, int32_t &offsetBias);
matmul::Matmul<matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, aType>,
matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, bType>,
matmul::MatmulType<AscendC::TPosition::VECIN, CubeFormat::ND, cType>,
matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, biasType>>
matmulObj;
AscendC::GlobalTensor<aType> aGlobal;
AscendC::GlobalTensor<bType> bGlobal;
AscendC::GlobalTensor<cType> cGlobal;
AscendC::GlobalTensor<biasType> biasGlobal;
AscendC::LocalTensor<cType> reluOutLocal;
TCubeTiling tiling;
AscendC::TQue<AscendC::TPosition::VECOUT, 1> reluOutQueue_;
};
/**
* @brief Set matmulLeaky input and output gm addr of current core.
* @param a: A matrix gm addr.
* @param b: B matrix gm addr.
* @param bias: Bias gm addr.
* @param c: C matrix gm addr.
* @param workspace: Temporary gm space addr required by matmul calc.
* @param tiling: matmul tiling data.
* @param pipe: Global memory and sync management TPipe object.
* @retval None
*/
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias,
GM_ADDR c, GM_ADDR workspace,
const TCubeTiling &tiling, AscendC::TPipe *pipe)
{
this->tiling = tiling;
aGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ aType *>(a), tiling.M * tiling.Ka);
bGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ bType *>(b), tiling.Kb * tiling.N);
cGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ cType *>(c), tiling.M * tiling.N);
biasGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ biasType *>(bias), tiling.N);
int32_t offsetA, offsetB, offsetC, offsetBias;
CalcOffset(AscendC::GetBlockIdx(), tiling, offsetA, offsetB, offsetC, offsetBias); // Calculate the gm offset based on the blockidx.
aGlobal = aGlobal[offsetA];
bGlobal = bGlobal[offsetB];
cGlobal = cGlobal[offsetC];
biasGlobal = biasGlobal[offsetBias];
pipe->InitBuffer(reluOutQueue_, 1, tiling.baseM * tiling.baseN * sizeof(cType)); // Init output buffer.
}
/**
* @brief Main process of matmul calculation
* @param pipe: Global memory and sync management TPipe object.
* @retval None
*/
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Process(AscendC::TPipe *pipe)
{
uint32_t computeRound = 0;
AscendC::printf("-----------%d\n", computeRound);
matmulObj.SetTensorA(aGlobal);
matmulObj.SetTensorB(bGlobal);
matmulObj.SetBias(biasGlobal);
while (matmulObj.template Iterate<true>()) { // Once Iterate, compute baseM * baseN, sync is set true here.
MatmulCompute(); // Get matmul compute result.
LeakyReluCompute(); // Compute leakyRelu.
CopyOut(computeRound); // Copy leakyRelu out result to GM.
computeRound++;
}
matmulObj.End();
}
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::MatmulCompute()
{
reluOutLocal = reluOutQueue_.AllocTensor<cType>();
matmulObj.template GetTensorC<true>(reluOutLocal, false, true);
}
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::LeakyReluCompute()
{
LeakyRelu(reluOutLocal, reluOutLocal, (cType)0.001, tiling.baseM * tiling.baseN);
reluOutQueue_.EnQue(reluOutLocal);
}
/**
* @brief Copy leakyRelu out result to GM.
* @param count: Iterate count(once Iterate, compute baseM * baseN).
* @retval None
*/
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::CopyOut(uint32_t count)
{
reluOutQueue_.DeQue<cType>();
const uint32_t roundM = tiling.singleCoreM / tiling.baseM;
const uint32_t roundN = tiling.singleCoreN / tiling.baseN;
uint32_t startOffset = (count % roundM * tiling.baseM * tiling.N + count / roundM * tiling.baseN);
AscendC::DataCopyParams copyParam = {(uint16_t)tiling.baseM, (uint16_t)(tiling.baseN * sizeof(cType) / AscendC::DEFAULT_C0_SIZE), 0,
(uint16_t)((tiling.N - tiling.baseN) * sizeof(cType) / AscendC::DEFAULT_C0_SIZE)};
DataCopy(cGlobal[startOffset], reluOutLocal, copyParam);
reluOutQueue_.FreeTensor(reluOutLocal);
}
/**
* @brief Calculate the gm offset based on the blockidx.
* @param blockIdx: Current Core blockidx.
* @param tiling: Matmul tiling data.
* @param offsetA: Gm offset of A matrix.
* @param offsetB: Gm offset of B matrix.
* @param offsetC: Gm offset of C matrix.
* @param offsetBias: Gm offset of Bias matrix.
* @retval None
*/
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void
MatmulLeakyKernel<aType, bType, cType, biasType>::CalcOffset(int32_t blockIdx, const TCubeTiling &tiling,
int32_t &offsetA, int32_t &offsetB, int32_t &offsetC,
int32_t &offsetBias)
{
auto mSingleBlocks = Ceiling(tiling.M, tiling.singleCoreM);
auto mCoreIndx = blockIdx % mSingleBlocks;
auto nCoreIndx = blockIdx / mSingleBlocks;
offsetA = mCoreIndx * tiling.Ka * tiling.singleCoreM;
offsetB = nCoreIndx * tiling.singleCoreN;
offsetC = mCoreIndx * tiling.N * tiling.singleCoreM + nCoreIndx * tiling.singleCoreN;
offsetBias = nCoreIndx * tiling.singleCoreN;
}
/**
* @brief matmul_leakyrelu kernel function entry
* @param a: A matrix gm addr.
* @param b: B matrix gm addr.
* @param bias: Bias gm addr.
* @param c: Out gm addr.
* @param workspace: Temporary gm space addr required by matmul calc.
* @param tiling: Tiling data.
* @retval None
*/
__global__ __mix__(1, 2) void matmul_leakyrelu_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c,
__kfc_workspace__ GM_ADDR workspace, AscendC::tiling::TCubeTiling tiling)
{
AscendC::TPipe pipe;
MatmulLeakyKernel<half, half, float, float> matmulLeakyKernel;
matmulLeakyKernel.Init(a, b, bias, c, workspace, tiling, &pipe);
REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulLeakyKernel.matmulObj, &matmulLeakyKernel.tiling); // Initialize the matmul object.
matmulLeakyKernel.Process(&pipe);
}
/**
* @brief Generate matmul tiling.
* @param ascendcPlatform: platform info.
*/
AscendC::tiling::TCubeTiling GenerateTiling(platform_ascendc::PlatformAscendC* ascendcPlatform)
{
using TPosition = matmul_tiling::TPosition;
using CubeFormat = matmul_tiling::CubeFormat;
using DataType = matmul_tiling::DataType;
int M = 1024;
int N = 640;
int K = 256;
TPosition leftPosition = TPosition::GM;
CubeFormat leftFormat = CubeFormat::ND;
DataType leftDtype = DataType::DT_FLOAT16;
bool isTransA = false;
TPosition rightPosition = TPosition::GM;
CubeFormat rightFormat = CubeFormat::ND;
DataType rightDtype = DataType::DT_FLOAT16;
bool isTransB = false;
TPosition resultPosition = TPosition::GM;
CubeFormat resultFormat = CubeFormat::ND;
DataType resultDtype = DataType::DT_FLOAT;
TPosition biasPosition = TPosition::GM;
CubeFormat biasFormat = CubeFormat::ND;
DataType biasDtype = DataType::DT_FLOAT;
bool isBias = true;
int usedCoreNum = 2;
int baseM = 256;
int baseN = 128;
matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform);
tilingApi.SetDim(usedCoreNum); // Set the number of cores that participate in multi-core computaion is 2.
tilingApi.SetAType(leftPosition, leftFormat, leftDtype, isTransA);
tilingApi.SetBType(rightPosition, rightFormat, rightDtype, isTransB);
tilingApi.SetCType(resultPosition, resultFormat, resultDtype);
tilingApi.SetBiasType(biasPosition, biasFormat, biasDtype);
tilingApi.SetOrgShape(M, N, K);
tilingApi.SetShape(M, N, K);
tilingApi.SetBias(isBias);
tilingApi.SetTraverse(matmul_tiling::MatrixTraverse::FIRSTM); // Set the matmul travse is FIRSTM.
tilingApi.SetFixSplit(baseM, baseN, -1); // Set the fixed baseM=128, baseN=256.
tilingApi.SetBufferSpace(-1, -1, -1);
AscendC::tiling::TCubeTiling tilingData;
int64_t res = tilingApi.GetTiling(tilingData); // Get matmul tiling data.
if (res == -1) {
std::cout << "gen tiling failed" << std::endl;
}
tilingData.stepM = 1; // Set the matmul tiling stepM=1.
tilingData.stepN = 1; // Set the matmul tiling stepN=1.
return tilingData;
}
int32_t main(int32_t argc, char *argv[])
{
const char *socVersion = "Ascend910B1";
auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(socVersion);
size_t aFileSize = 262144 * sizeof(int16_t);
size_t bFileSize = 163840 * sizeof(int16_t);
size_t cFileSize = 655360 * sizeof(float);
size_t biasFileSize = 640 * sizeof(float);
size_t tilingFileSize = sizeof(TCubeTiling);
size_t userWorkspaceSize = 0;
size_t systemWorkspaceSize = static_cast<size_t>(ascendcPlatform->GetLibApiWorkSpaceSize());
size_t workspaceSize = userWorkspaceSize + systemWorkspaceSize;
auto tiling = GenerateTiling(ascendcPlatform);
uint32_t numBlocks = 1;
aclInit(nullptr);
int32_t deviceId = 0;
aclrtSetDevice(deviceId);
aclrtStream stream = nullptr;
aclrtCreateStream(&stream);
uint8_t *inputAHost;
uint8_t *inputADevice;
aclrtMallocHost((void **)(&inputAHost), aFileSize);
aclrtMalloc((void **)&inputADevice, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
ReadFile("./input/x1_gm.bin", aFileSize, inputAHost, aFileSize);
aclrtMemcpy(inputADevice, aFileSize, inputAHost, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
uint8_t *inputBHost;
uint8_t *inputBDevice;
aclrtMallocHost((void **)(&inputBHost), bFileSize);
aclrtMalloc((void **)&inputBDevice, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
ReadFile("./input/x2_gm.bin", bFileSize, inputBHost, bFileSize);
aclrtMemcpy(inputBDevice, bFileSize, inputBHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
uint8_t *outputCHost;
uint8_t *outputCDevice;
aclrtMallocHost((void **)(&outputCHost), cFileSize);
aclrtMalloc((void **)&outputCDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
uint8_t *inputBiasHost;
uint8_t *inputBiasDevice;
aclrtMallocHost((void **)(&inputBiasHost), biasFileSize);
aclrtMalloc((void **)&inputBiasDevice, biasFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
ReadFile("./input/bias.bin", biasFileSize, inputBiasHost, biasFileSize);
aclrtMemcpy(inputBiasDevice, biasFileSize, inputBiasHost, biasFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
uint8_t *workspaceDevice;
aclrtMalloc((void **)&workspaceDevice, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
matmul_leakyrelu_custom<<<numBlocks, nullptr, stream>>>(inputADevice, inputBDevice, inputBiasDevice, outputCDevice,
workspaceDevice, tiling);
aclrtSynchronizeStream(stream);
aclrtFree(inputADevice);
aclrtFreeHost(inputAHost);
aclrtFree(inputBDevice);
aclrtFreeHost(inputBHost);
aclrtMemcpy(outputCHost, cFileSize, outputCDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
WriteFile("./output/output.bin", outputCHost, cFileSize);
aclrtFree(outputCDevice);
aclrtFreeHost(outputCHost);
aclrtFree(inputBiasDevice);
aclrtFreeHost(inputBiasHost);
aclrtFree(workspaceDevice);
aclrtDestroyStream(stream);
aclrtResetDevice(deviceId);
aclFinalize();
return 0;
}
