--- title: WelfordFinalize Tiling description: "Ascend C提供WelfordFinalize Tiling API，方便用户获取WelfordFinalize kernel计算时所需的Tiling参数。" url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/API/ascendcopapi/atlasascendc_api_07_0815.html sourcePath: /source/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0815.html indexId: c2ac79eb0c894142fb224043e783ccca31817c5da3ffa99fe2199d858c4f1c8e76 --- # WelfordFinalize Tiling #### 功能说明 Ascend C提供WelfordFinalize Tiling API，方便用户获取WelfordFinalize kernel计算时所需的Tiling参数。 获取Tiling参数主要步骤如下： 具体为，通过GetWelfordFinalizeMaxMinTmpSize获取WelfordFinalize接口计算所需最大和最小临时空间大小。 kernel侧WelfordFinalize接口的计算需要开发者预留/申请临时空间，GetWelfordFinalizeMaxMinTmpSize用于在host侧获取预留/申请的最大最小临时空间大小，开发者基于此范围选择合适的空间大小作为Tiling参数传递到kernel侧使用。- 为保证功能正确，预留/申请的临时空间大小不能小于最小临时空间大小； - 在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。 #### 函数原型 ``` void GetWelfordFinalizeMaxMinTmpSize(const ge::Shape& srcShape, const uint32_t typeSize, const bool isReuseSource, uint32_t& maxValue, uint32_t& minValue) ``` #### 参数说明 **表1 GetWelfordFinalizeMaxMinTmpSize接口参数说明** | 参数名 | 输入/输出 | 描述 | | --- | --- | --- | | srcShape | 输入 | 输入inputMean/inputVariance的shape信息{abLength}。 | | typeSize | 输入 | 输入inputMean/inputVariance的数据类型大小，单位为字节。比如输入的数据类型为float，此处应传入4。 | | isReuseSource | 输入 | 是否允许修改源操作数。该参数取值与WelfordFinalize接口一致。 | | maxValue | 输出 | WelfordFinalize接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间\-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。最大空间大小为0表示计算不需要临时空间。 说明： maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 | | minValue | 输出 | WelfordFinalize接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。最小空间大小为0表示计算不需要临时空间。 | #### 返回值说明 无 #### 约束说明 无 #### 调用示例 1. 将WelfordFinalizeTiling结构体参数增加至TilingData结构体，作为TilingData结构体的一个字段。 1 2 3 4 5 6 7 8 9 10 11 BEGIN_TILING_DATA_DEF(WelfordFinalizeCustomTilingData) // 注册一个tiling的类，以tiling的名字作为入参 TILING_DATA_FIELD_DEF(uint32_t, isCounts); // 添加tiling字段 TILING_DATA_FIELD_DEF(uint32_t, rnLength); TILING_DATA_FIELD_DEF(uint32_t, abLength); TILING_DATA_FIELD_DEF(uint32_t, rLength); TILING_DATA_FIELD_DEF(uint32_t, head); TILING_DATA_FIELD_DEF(uint32_t, headLength); TILING_DATA_FIELD_DEF(uint32_t, tail); TILING_DATA_FIELD_DEF(uint32_t, tailLength); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(WelfordFinalizeCustom, WelfordFinalizeCustomTilingData)// 将WelfordFinalizeCustomTilingData结构体参数增加至TilingData结构体 2. Tiling实现函数中，首先调用GetWelfordFinalizeMaxMinTmpSize接口获取WelfordFinalize接口能完成计算所需最大/最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小，然后根据输入shape、剩余的可供计算的空间大小等信息获取WelfordFinalize kernel侧接口所需tiling参数。 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext *context) { WelfordFinalizeCustomTilingData tiling; const gert::RuntimeAttrs *attrs = context->GetAttrs(); const uint32_t isCounts = *(attrs->GetAttrPointer(0)); const uint32_t rnLength = *(attrs->GetAttrPointer(1)); const uint32_t abLength = *(attrs->GetAttrPointer(2)); const uint32_t rLength = *(attrs->GetAttrPointer(3)); const uint32_t head = *(attrs->GetAttrPointer(4)); const uint32_t headLength = *(attrs->GetAttrPointer(5)); const uint32_t tail = *(attrs->GetAttrPointer(6)); const uint32_t tailLength = *(attrs->GetAttrPointer(7)); std::vector srcDims = {abLength}; ge::Shape srcShape(srcDims); // 本样例中仅作为样例说明，通过GetWelfordFinalizeMaxMinTmpSize获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小 uint32_t maxTmpsize = 0; uint32_t minTmpsize = 0; AscendC::GetWelfordFinalizeMaxMinTmpSize(srcShape, 4, false, maxTmpsize, minTmpsize); // auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo()); // AscendC::GetWelfordFinalizeMaxMinTmpSize(srcShape, 4, false, ascendcPlatform, maxTmpsize, minTmpsize); ... // 其他逻辑 context->SetTilingKey(1); tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); size_t *currentWorkspace = context->GetWorkspaceSizes(1); currentWorkspace[0] = 0; return ge::GRAPH_SUCCESS; } } // namespace optiling 3. 对应的kernel侧通过在核函数中调用GET_TILING_DATA获取TilingData，继而将TilingData中的WelfordFinalize Tiling信息传入WelfordFinalize接口参与计算。完整的kernel侧样例请参考WelfordFinalize。 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 extern "C" __global__ __aicore__ void welford_finalize_custom( GM_ADDR inputX_gm, GM_ADDR mean_gm, GM_ADDR var_gm, GM_ADDR outputMean_gm, GM_ADDR outputVariance_gm, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); if (TILING_KEY_IS(1)) { if (tilingData.isCounts) { KernelWelfordFinalize op; op.Init(inputX_gm, mean_gm, var_gm, outputMean_gm, outputVariance_gm, tilingData.rnLength, tilingData.abLength, tilingData.rLength, tilingData.head, tilingData.headLength, tilingData.tail, tilingData.tailLength); op.Process(); } else { KernelWelfordFinalize op; op.Init(inputX_gm, mean_gm, var_gm, outputMean_gm, outputVariance_gm, tilingData.rnLength, tilingData.abLength, tilingData.rLength, tilingData.head, tilingData.headLength, tilingData.tail, tilingData.tailLength); op.Process(); } } }