Ascend C提供一组LayerNorm Tiling API,方便用户获取LayerNorm kernel计算时所需的Tiling参数。
获取Tiling参数主要分为如下两步:
LayerNorm Tiling结构体的定义如下,开发者无需关注该Tiling结构的具体信息,只需要传递到kernel侧,传入LayerNorm高阶API接口,直接进行使用即可。
struct LayerNormTiling { uint32_t bLength = 0; uint32_t sLength = 0; uint32_t hLength = 0; uint32_t originalHLength = 0; uint32_t inputXSize = 0; uint32_t meanVarSize = 0; uint32_t numberOfTmpBuf = 0; uint32_t meanTmpTensorPos = 0; uint32_t meanTmpTensorSize = 0; uint32_t varianceTmpTensorPos = 0; uint32_t varianceTmpTensorSize = 0; uint32_t tmpBufSize = 0; uint32_t oneTmpSize = 0; uint32_t firstTmpStartPos = 0; uint32_t secondTmpStartPos = 0; uint32_t thirdTmpStartPos = 0; uint32_t loopRound = 0; uint32_t inputRoundSize = 0; uint32_t inputTailSize = 0; uint32_t inputTailPos = 0; uint32_t meanVarRoundSize = 0; uint32_t meanVarTailSize = 0; uint32_t meanVarTailPos = 0; uint32_t bshCurLength = 0; uint32_t bsCurLength = 0; uint32_t lastDimValueBack = 0; };
1 | void GetLayerNormMaxMinTmpSize(const ge::Shape& srcShape, const uint32_t typeSize, const bool isReuseSource, uint32_t& maxValue, uint32_t& minValue) |
1 | void GetLayerNormNDTillingInfo(const ge::Shape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, const bool isReuseSource, optiling::LayerNormTiling& tilling) |
接口 |
输入/输出 |
功能 |
---|---|---|
srcShape |
输入 |
输入数据inputX的shape信息{B, S, storageHLength, originHLength},包括当前输入的inputX的shape信息,以及地址对齐前(如存在H轴补齐操作)的原有shape信息 。 在API支持的场景下,storageHLength和originHLength保持一致。 |
typeSize |
输入 |
算子输入的数据类型大小,单位为字节。比如算子输入的数据类型为half,此处应传入2。 |
isReuseSource |
输入 |
是否复用源操作数的内存空间,与LayerNorm接口一致。 |
maxValue |
输出 |
输出LayerNorm接口所需的tiling信息(最大临时空间大小)。 LayerNorm接口能完成计算所需的最大临时空间大小,超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内,随着临时空间增大,kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能,开发者可以根据实际的内存使用情况进行空间预留/申请。 说明:
maxValue仅作为参考值,有可能大于Unified Buffer剩余空间的大小,该场景下,开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 |
minValue |
输出 |
输出LayerNorm接口所需的tiling信息(最小临时空间大小)。 LayerNorm接口能完成计算所需最小临时空间大小。为保证功能正确,接口计算时预留/申请的临时空间不能小于该数值。 |
参数名称 |
输入/输出 |
含义 |
---|---|---|
srcShape |
输入 |
输入数据inputX的shape信息{B, S, storageHLength, originHLength},包括当前输入的inputX的shape信息,以及地址对齐前(如存在H轴补齐操作)的原有shape信息。 |
stackBufferSize |
输入 |
可供Layernorm接口使用的空间大小,单位Byte。 |
typeSize |
输入 |
算子输入的数据类型大小,单位为字节。比如算子输入的数据类型为half,此处应传入2。 |
isReuseSource |
输入 |
是否可以复用inputX的内存空间。 |
tilling |
输出 |
输入数据的切分信息。 |
无
如下样例介绍了使用LayerNorm高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。样例中输入Tensor的shape大小为[2, 16, 64],输入的数据类型为half。
BEGIN_TILING_DATA_DEF(TilingData) // 注册一个tiling的类,以tiling的名字作为入参 TILING_DATA_FIELD_DEF(uint32_t, totalLength); // 添加tiling字段,总计算数据量 TILING_DATA_FIELD_DEF(uint32_t, tileNum); // 添加tiling字段,每个核上总计算数据分块个数 ... // 添加其他tiling字段 TILING_DATA_FIELD_DEF_STRUCT(LayerNormTiling, layernormTilingData); // 将LayerNormTiling结构体参数增加至TilingData结构体 END_TILING_DATA_DEF;
namespace optiling { const uint32_t BLOCK_DIM = 8; const uint32_t TILE_NUM = 8; static ge::graphStatus TilingFunc(gert::TilingContext* context) { TilingData tiling; uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize(); context->SetBlockDim(BLOCK_DIM); tiling.set_totalLength(totalLength); tiling.set_tileNum(TILE_NUM); // 设置其他Tiling参数 ... // {B, S, storageHLength, originHLength} std::vector<int64_t> shapeVec = {2, 16, 64, 64}; ge::Shape srcShape(shapeVec); // 本样例中仅做为样例说明,通过GetLayerNormMaxMinTmpSize获取最小值并传入,来保证功能正确,开发者可以根据需要传入合适的空间大小 uint32_t max; uint32_t min; AscendC::GetLayerNormMaxMinTmpSize(srcShape, sizeof(half), false, max, min); // 获取Layernorm Tiling参数 AscendC::GetLayerNormNDTillingInfo(srcShape, min, sizeof(half), false, tiling.layernormTilingData); ... // 其他逻辑 tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); context->SetTilingKey(1); return ge::GRAPH_SUCCESS; } } // namespace optiling
extern "C" __global__ __aicore__ void func_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelFunc op; op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum,tilingData.layernormTilingData); if (TILING_KEY_IS(1)) { op.Process(); } }