CANN/asc-devkit: REGISTER_NONE_TILING API

REGISTER_NONE_TILING【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品xAtlas 推理系列产品 AI CorexAtlas 推理系列产品 Vector CorexAtlas 训练系列产品x功能说明在Kernel侧使用标准C语法自定义的TilingData结构体时若用户不确定需要注册哪些结构体可使用该接口告知框架侧需使用未注册的标准C语法来定义TilingData并配套GET_TILING_DATA_WITH_STRUCTGET_TILING_DATA_MEMBERGET_TILING_DATA_PTR_WITH_STRUCT来获取对应的TilingData。函数原型REGISTER_NONE_TILING参数说明无约束说明暂不支持Kernel直调工程。使用GET_TILING_DATA需提供默认注册的TilingData结构体但本接口不注册TilingData结构体故不支持与5.11.1-GET_TILING_DATA组合使用。不支持和REGISTER_TILING_DEFAULT或REGISTER_TILING_FOR_TILINGKEY混用即不支持注册TilingData结构体的场景与非注册场景混合使用。调用示例# Tiling模板库提供方无法预知用户实例化何种TilingData结构体 template class BrcDag struct BroadcastBaseTilingData { int32_t scheMode; int32_t shapeLen; int32_t ubSplitAxis; int32_t ubFormer; int32_t ubTail; int64_t ubOuter; int64_t blockFormer; int64_t blockTail; int64_t dimProductBeforeUbInner; int64_t elemNum; int64_t blockNum; int64_t outputDims[BROADCAST_MAX_DIMS_NUM]; int64_t outputStrides[BROADCAST_MAX_DIMS_NUM]; int64_t inputDims[BrcDag::InputSize][2]; // 整块 尾块 int64_t inputBrcDims[BrcDag::CopyBrcSize][BROADCAST_MAX_DIMS_NUM]; int64_t inputVecBrcDims[BrcDag::VecBrcSize][BROADCAST_MAX_DIMS_NUM]; int64_t inputStrides[BrcDag::InputSize][BROADCAST_MAX_DIMS_NUM]; int64_t inputBrcStrides[BrcDag::CopyBrcSize][BROADCAST_MAX_DIMS_NUM]; int64_t inputVecBrcStrides[BrcDag::VecBrcSize]; char scalarData[BROADCAST_MAX_SCALAR_BYTES]; }; template uint64_t schMode, class BrcDag class BroadcastSch { public: __aicore__ inline explicit BroadcastSch(GM_ADDR tmpTiling) : tiling(tmpTiling) {} template class... Args __aicore__ inline void Process(Args... args) { REGISTER_NONE_TILING; // 告知框架侧使用未注册的TilingData结构体 if constexpr (schMode 1) { GET_TILING_DATA_WITH_STRUCT(BroadcastBaseTilingDataBrcDag, tilingData, tiling); GET_TILING_DATA_MEMBER(BroadcastBaseTilingDataBrcDag, blockNum, blockNumVar, tiling); TPipe pipe; BroadcastNddmaSchBrcDag, false sch(tilingData); // 获取Schedule sch.Init(pipe, args...); sch.Process(); } else if constexpr (schMode 202) { GET_TILING_DATA_PTR_WITH_STRUCT(BroadcastOneDimTilingDataAdvance, tilingDataPtr, tiling); BroadcastOneDimAdvanceSchBrcDag sch(tilingDataPtr); // 获取Schedule sch.Init(args...); sch.Process(); } } public: GM_ADDR tiling; };#用户通过传入schMode, OpDag模板参数来实例化模板库 using namespace AscendC; template uint64_t schMode __global__ __aicore__ void mul(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { if constexpr (std::is_sameDTYPE_X1, int8_t::value) { // int8 using OpDag MulDag::MulInt8Op::OpDag; BroadcastSchschMode, OpDag sch(tiling); sch.Process(x1, x2, y); } else if constexpr (std::is_sameDTYPE_X1, uint8_t::value) { // uint8 using OpDag MulDag::MulUint8Op::OpDag; BroadcastSchschMode, OpDag sch(tiling); sch.Process(x1, x2, y); } }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考