TopK Tiling【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit功能说明用于获取TopK Tiling参数。Ascend C提供TopK Tiling API方便用户获取TopK kernel计算时所需的Tiling参数。获取Tiling参数主要分为如下两步获取TopK接口计算所需最小和最大临时空间大小注意该步骤不是必须的只是作为一个参考供合理分配计算空间。获取TopK kernel侧接口所需tiling参数。TopK Tiling结构体的定义如下开发者无需关注该tiling结构的具体信息只需要传递到kernel侧传入TopK高阶API接口直接进行使用即可。struct TopkTiling { int32_t tmpLocalSize 0; int32_t allDataSize 0; int32_t innerDataSize 0; uint32_t sortRepeat 0; int32_t mrgSortRepeat 0; int32_t kAlignFourBytes 0; int32_t kAlignTwoBytes 0; int32_t maskOffset 0; int32_t maskVreducev2FourBytes 0; int32_t maskVreducev2TwoBytes 0; int32_t mrgSortSrc1offset 0; int32_t mrgSortSrc2offset 0; int32_t mrgSortSrc3offset 0; int32_t mrgSortTwoQueueSrc1Offset 0; int32_t mrgFourQueueTailPara1 0; int32_t mrgFourQueueTailPara2 0; int32_t srcIndexOffset 0; uint32_t copyUbToUbBlockCount 0; int32_t topkMrgSrc1MaskSizeOffset 0; int32_t topkNSmallSrcIndexOffset 0; uint32_t vreduceValMask0 0; uint32_t vreduceValMask1 0; uint32_t vreduceIdxMask0 0; uint32_t vreduceIdxMask1 0; uint16_t vreducehalfValMask0 0; uint16_t vreducehalfValMask1 0; uint16_t vreducehalfValMask2 0; uint16_t vreducehalfValMask3 0; uint16_t vreducehalfValMask4 0; uint16_t vreducehalfValMask5 0; uint16_t vreducehalfValMask6 0; uint16_t vreducehalfValMask7 0; };函数原型bool GetTopKMaxMinTmpSize(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const bool isReuseSource, const bool isInitIndex, enum TopKMode mode, const bool isLargest, const uint32_t dataTypeSize, uint32_t maxValue, uint32_t minValue)bool GetTopKMaxMinTmpSize(const int32_t inner, const int32_t outter, const int32_t k, const bool isReuseSource, const bool isInitIndex, enum TopKMode mode, const bool isLargest, ge::DataType dataType, const TopKConfig config, uint32_t maxValue, uint32_t minValue)bool TopKTilingFunc(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const int32_t k, const uint32_t dataTypeSize, const bool isInitIndex, enum TopKMode mode, const bool isLargest, optiling::TopkTiling topKTiling)bool TopKTilingFunc(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const int32_t k, const uint32_t dataTypeSize, const bool isInitIndex, enum TopKMode mode, const bool isLargest, AscendC::tiling::TopkTiling topKTiling)参数说明表1GetTopKMaxMinTmpSize接口参数列表接口输入/输出功能ascendcPlatform输入传入硬件平台的信息PlatformAscendC定义请参见构造及析构函数。inner输入表示TopK接口输入srcLocal的内轴长度该参数的取值为32的整数倍。outter输入表示TopK接口输入srcLocal的外轴长度。k输入获取前k个最大值或最小值及其对应的索引。isReuseSource输入中间变量是否能够复用输入内存。与kernel侧接口的isReuseSrc保持一致。isInitIndex输入是否传入输入数据对应的索引与kernel侧接口一致。mode输入选择TopKMode::TOPK_NORMAL模式或者TopKMode::TOPK_NSMALL模式与kernel侧接口一致。isLargest输入表示降序/升序true表示降序false表示升序。与kernel侧接口一致。dataType输入表示待排序数据的数据类型。该参数的取值与Kernel接口参数srcLocal的数据类型保持一致。config输入TopK计算的相关配置TopKConfig类型定义如下方代码所示包括算法选择、取最大值或最小值、是否对结果排序。该参数的配置需要与TopK Kernel接口模板参数的配置保持一致。algo选择的排序算法。默认为MERGE_SORT算法当前仅支持RADIX_SELECT算法用户需要显式指定algo为TopKAlgo::RADIX_SELECT。order表示获取前k个最大值或者获取前k个最小值取值如下UNSET默认值按照函数参数isLargest的配置实现。isLargest为true时取前k个最大值及其对应的索引isLargest为false取前k个最小值及其对应的索引。LARGEST表示取前k个最大值及其对应的索引。取值为LARGEST时函数参数isLargest的配置不生效。SMALLEST表示取前k个最小值及其对应的索引。取值为SMALLEST时函数参数isLargest的配置不生效。sorted表示是否对输出结果进行排序。取值为true对输出结果进行排序取值为false不对输出结果进行排序。dataTypeSize输入参与计算的srcLocal数据类型的大小比如half2 float4maxValue输出TopK接口内部完成计算需要的最大临时空间大小单位是Byte。说明maxValue仅作为参考值有可能大于Unified Buffer剩余空间的大小该场景下开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。minValue输出TopK接口内部完成计算需要的最小临时空间大小单位是Byte。struct TopKConfig { TopKAlgo algo TopKAlgo::MERGE_SORT; TopKOrder order TopKOrder::UNSET; bool sorted true; }; enum class TopKAlgo { RADIX_SELECT, MERGE_SORT }; enum class TopKOrder { UNSET, LARGEST, SMALLEST };表2TopKTilingFunc接口参数列表接口输入/输出功能ascendcPlatform输入传入硬件平台的信息PlatformAscendC定义请参见构造及析构函数。inner输入表示TopK接口输入srcLocal的内轴长度该参数的取值为32的整数倍。outter输入表示TopK接口输入srcLocal的外轴长度。k输入获取前k个最大值或最小值及其对应的索引。dataTypeSize输入参与计算的srcLocal数据类型的大小比如half2 float4。isInitIndex输入是否传入输入数据对应的索引与kernel侧接口一致。mode输入选择TopKMode::TOPK_NORMAL模式或者TopKMode::TOPK_NSMALL模式与kernel侧接口一致。isLargest输入表示降序/升序true表示降序false表示升序。与kernel侧接口一致。topKTiling输出输出TopK接口所需的tiling信息。返回值说明GetTopKMaxMinTmpSize返回值为true/falsetrue表示成功拿到TopK接口内部计算需要的最大和最小临时空间大小false表示获取失败。TopKTilingFunc返回值为true/falsetrue表示成功拿到TopK的Tiling各项参数值false表示获取失败。约束说明无调用示例如下样例介绍了使用TopK高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。将TopK Tiling结构体参数增加至TilingData结构体作为TilingData结构体的一个字段。namespace optiling { BEGIN_TILING_DATA_DEF(TilingData) TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, tilenum); // 添加其他tiling字段 ... TILING_DATA_FIELD_DEF(int32_t, k); TILING_DATA_FIELD_DEF(bool, islargest); TILING_DATA_FIELD_DEF(bool, isinitindex); TILING_DATA_FIELD_DEF(bool, ishasfinish); TILING_DATA_FIELD_DEF(uint32_t, tmpsize); TILING_DATA_FIELD_DEF(int32_t, outter); TILING_DATA_FIELD_DEF(int32_t, inner); TILING_DATA_FIELD_DEF(int32_t, n); TILING_DATA_FIELD_DEF(int32_t, order); TILING_DATA_FIELD_DEF(int32_t, sorted); TILING_DATA_FIELD_DEF_STRUCT(TopkTiling, topkTilingData); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(TopkCustom, TilingData) } // namespace optilingTiling实现函数中首先调用GetTopKMaxMinTmpSize接口获取TopK接口能完成计算所需最大/最小临时空间大小根据该范围结合实际的内存使用情况设置合适的空间大小然后根据输入shape等信息获取TopK kernel侧接口所需tiling参数。MERGE_SORT算法参考如下调用示例。namespace optiling { const uint32_t NUM_BLOCKS 8; const uint32_t TILE_NUM 8; const int32_t OUTTER 2; const int32_t INNER 32; const int32_t N 32; const int32_t K 8; const bool IS_LARGEST true; const bool IS_INITINDEX true; const bool IS_REUSESOURCE false; static ge::graphStatus TilingFunc(gert::TilingContext* context) { TilingData tiling; uint32_t totalLength context-GetInputTensor(0)-GetShapeSize(); context-SetBlockDim(NUM_BLOCKS); tiling.set_totalLength(totalLength); tiling.set_tileNum(TILE_NUM); tiling.set_k(K); tiling.set_outter(OUTTER); tiling.set_inner(INNER); tiling.set_n(N); tiling.set_islargest(IS_LARGEST); tiling.set_isinitindex(IS_INITINDEX); // 设置其他Tiling参数 ... // 本样例中仅作为样例说明通过GetTopKMaxMinTmpSize获取最小值并传入来保证功能正确开发者可以根据需要传入合适的空间大小。 uint32_t maxsize 0; uint32_t minsize 0; uint32_t dtypesize 4; // float类型 auto ascendcPlatform platform_ascendc::PlatformAscendC(context-GetPlatformInfo()); AscendC::TopKTilingFunc( ascendcPlatform, tiling.inner, tiling.outter, tiling.k, dtypesize, tiling.isinitindex, AscendC::TopKMode::TOPK_NSMALL, tiling.islargest, tiling.topkTilingData); AscendC::GetTopKMaxMinTmpSize( ascendcPlatform, tiling.inner, tiling.outter, IS_REUSESOURCE, tiling.isinitindex, AscendC::TopKMode::TOPK_NSMALL, tiling.islargest, dtypesize, maxsize, minsize); tiling.set_tmpsize(minsize); ... // 其他逻辑 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 optilingRADIX_SELECT算法参考如下调用示例。namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext* context) { std::mapge::DataType, uint32_t dtypeSizes {{ge::DataType::DT_UINT32, 4}, {ge::DataType::DT_INT32, 4}}; RadixtopkCustomTilingData tiling; const gert::RuntimeAttrs* attrs context-GetAttrs(); const uint32_t is_init_index *(attrs-GetAttrPointeruint32_t(0)); const uint32_t is_reuse_src *(attrs-GetAttrPointeruint32_t(1)); const uint32_t order *(attrs-GetAttrPointeruint32_t(2)); const uint32_t is_largest *(attrs-GetAttrPointeruint32_t(3)); const uint32_t outter *(attrs-GetAttrPointeruint32_t(4)); const uint32_t inner *(attrs-GetAttrPointeruint32_t(5)); const uint32_t n *(attrs-GetAttrPointeruint32_t(6)); const uint32_t k *(attrs-GetAttrPointeruint32_t(7)); const uint32_t k_pad *(attrs-GetAttrPointeruint32_t(8)); const uint32_t sorted *(attrs-GetAttrPointeruint32_t(9)); const uint32_t top_mode *(attrs-GetAttrPointeruint32_t(10)); auto xDType context-GetInputTensor(0)-GetDataType(); uint32_t typeSize dtypeSizes.at(xDType); AscendC::TopKConfig config; config.algo AscendC::TopKAlgo::RADIX_SELECT; if (order 1) { config.order AscendC::TopKOrder::LARGEST; } else if (order 2) { config.order AscendC::TopKOrder::SMALLEST; } else { config.order AscendC::TopKOrder::UNSET; } if (sorted 0) { config.sorted false; } else { config.sorted true; } uint32_t maxValue 0; uint32_t minValue 0; if (top_mode 0) { AscendC::GetTopKMaxMinTmpSize( inner, outter, k, is_reuse_src, is_init_index, AscendC::TopKMode::TOPK_NORMAL, is_largest, xDType, config, maxValue, minValue); context-SetTilingKey(0); } else { AscendC::GetTopKMaxMinTmpSize( inner, outter, k, is_reuse_src, is_init_index, AscendC::TopKMode::TOPK_NSMALL, is_largest, xDType, config, maxValue, minValue); context-SetTilingKey(1); } context-SetBlockDim(1); tiling.set_is_init_index(is_init_index); tiling.set_is_reuse_src(is_reuse_src); tiling.set_order(order); tiling.set_is_largest(is_largest); tiling.set_outter(outter); tiling.set_inner(inner); tiling.set_n(n); tiling.set_k(k); tiling.set_k_pad(k_pad); tiling.set_sorted(sorted); tiling.set_top_mode(top_mode); 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对应的kernel侧通过在核函数中调用GET_TILING_DATA获取TilingData继而将TilingData中的TopK Tiling信息传入TopK接口参与计算。完整的kernel侧样例请参考调用示例。extern C __global__ __aicore__ void topk_custom( GM_ADDR srcVal, GM_ADDR srcIdx, GM_ADDR finishLocal, GM_ADDR dstVal, GM_ADDR dstIdx, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelTopKfloat, true, true, false, false, AscendC::TopKMode::TOPK_NSMALL op; op.Init( srcVal, srcIdx, finishLocal, dstVal, dstIdx, tilingData.k, tilingData.islargest, tilingData.tmpsize, tilingData.outter, tilingData.inner, tilingData.n, tilingData.topkTilingData); op.Process(); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN Ascend C TopK Tiling参数
发布时间：2026/7/17 8:50:54
TopK Tiling【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit功能说明用于获取TopK Tiling参数。Ascend C提供TopK Tiling API方便用户获取TopK kernel计算时所需的Tiling参数。获取Tiling参数主要分为如下两步获取TopK接口计算所需最小和最大临时空间大小注意该步骤不是必须的只是作为一个参考供合理分配计算空间。获取TopK kernel侧接口所需tiling参数。TopK Tiling结构体的定义如下开发者无需关注该tiling结构的具体信息只需要传递到kernel侧传入TopK高阶API接口直接进行使用即可。struct TopkTiling { int32_t tmpLocalSize 0; int32_t allDataSize 0; int32_t innerDataSize 0; uint32_t sortRepeat 0; int32_t mrgSortRepeat 0; int32_t kAlignFourBytes 0; int32_t kAlignTwoBytes 0; int32_t maskOffset 0; int32_t maskVreducev2FourBytes 0; int32_t maskVreducev2TwoBytes 0; int32_t mrgSortSrc1offset 0; int32_t mrgSortSrc2offset 0; int32_t mrgSortSrc3offset 0; int32_t mrgSortTwoQueueSrc1Offset 0; int32_t mrgFourQueueTailPara1 0; int32_t mrgFourQueueTailPara2 0; int32_t srcIndexOffset 0; uint32_t copyUbToUbBlockCount 0; int32_t topkMrgSrc1MaskSizeOffset 0; int32_t topkNSmallSrcIndexOffset 0; uint32_t vreduceValMask0 0; uint32_t vreduceValMask1 0; uint32_t vreduceIdxMask0 0; uint32_t vreduceIdxMask1 0; uint16_t vreducehalfValMask0 0; uint16_t vreducehalfValMask1 0; uint16_t vreducehalfValMask2 0; uint16_t vreducehalfValMask3 0; uint16_t vreducehalfValMask4 0; uint16_t vreducehalfValMask5 0; uint16_t vreducehalfValMask6 0; uint16_t vreducehalfValMask7 0; };函数原型bool GetTopKMaxMinTmpSize(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const bool isReuseSource, const bool isInitIndex, enum TopKMode mode, const bool isLargest, const uint32_t dataTypeSize, uint32_t maxValue, uint32_t minValue)bool GetTopKMaxMinTmpSize(const int32_t inner, const int32_t outter, const int32_t k, const bool isReuseSource, const bool isInitIndex, enum TopKMode mode, const bool isLargest, ge::DataType dataType, const TopKConfig config, uint32_t maxValue, uint32_t minValue)bool TopKTilingFunc(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const int32_t k, const uint32_t dataTypeSize, const bool isInitIndex, enum TopKMode mode, const bool isLargest, optiling::TopkTiling topKTiling)bool TopKTilingFunc(const platform_ascendc::PlatformAscendC ascendcPlatform, const int32_t inner, const int32_t outter, const int32_t k, const uint32_t dataTypeSize, const bool isInitIndex, enum TopKMode mode, const bool isLargest, AscendC::tiling::TopkTiling topKTiling)参数说明表1GetTopKMaxMinTmpSize接口参数列表接口输入/输出功能ascendcPlatform输入传入硬件平台的信息PlatformAscendC定义请参见构造及析构函数。inner输入表示TopK接口输入srcLocal的内轴长度该参数的取值为32的整数倍。outter输入表示TopK接口输入srcLocal的外轴长度。k输入获取前k个最大值或最小值及其对应的索引。isReuseSource输入中间变量是否能够复用输入内存。与kernel侧接口的isReuseSrc保持一致。isInitIndex输入是否传入输入数据对应的索引与kernel侧接口一致。mode输入选择TopKMode::TOPK_NORMAL模式或者TopKMode::TOPK_NSMALL模式与kernel侧接口一致。isLargest输入表示降序/升序true表示降序false表示升序。与kernel侧接口一致。dataType输入表示待排序数据的数据类型。该参数的取值与Kernel接口参数srcLocal的数据类型保持一致。config输入TopK计算的相关配置TopKConfig类型定义如下方代码所示包括算法选择、取最大值或最小值、是否对结果排序。该参数的配置需要与TopK Kernel接口模板参数的配置保持一致。algo选择的排序算法。默认为MERGE_SORT算法当前仅支持RADIX_SELECT算法用户需要显式指定algo为TopKAlgo::RADIX_SELECT。order表示获取前k个最大值或者获取前k个最小值取值如下UNSET默认值按照函数参数isLargest的配置实现。isLargest为true时取前k个最大值及其对应的索引isLargest为false取前k个最小值及其对应的索引。LARGEST表示取前k个最大值及其对应的索引。取值为LARGEST时函数参数isLargest的配置不生效。SMALLEST表示取前k个最小值及其对应的索引。取值为SMALLEST时函数参数isLargest的配置不生效。sorted表示是否对输出结果进行排序。取值为true对输出结果进行排序取值为false不对输出结果进行排序。dataTypeSize输入参与计算的srcLocal数据类型的大小比如half2 float4maxValue输出TopK接口内部完成计算需要的最大临时空间大小单位是Byte。说明maxValue仅作为参考值有可能大于Unified Buffer剩余空间的大小该场景下开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。minValue输出TopK接口内部完成计算需要的最小临时空间大小单位是Byte。struct TopKConfig { TopKAlgo algo TopKAlgo::MERGE_SORT; TopKOrder order TopKOrder::UNSET; bool sorted true; }; enum class TopKAlgo { RADIX_SELECT, MERGE_SORT }; enum class TopKOrder { UNSET, LARGEST, SMALLEST };表2TopKTilingFunc接口参数列表接口输入/输出功能ascendcPlatform输入传入硬件平台的信息PlatformAscendC定义请参见构造及析构函数。inner输入表示TopK接口输入srcLocal的内轴长度该参数的取值为32的整数倍。outter输入表示TopK接口输入srcLocal的外轴长度。k输入获取前k个最大值或最小值及其对应的索引。dataTypeSize输入参与计算的srcLocal数据类型的大小比如half2 float4。isInitIndex输入是否传入输入数据对应的索引与kernel侧接口一致。mode输入选择TopKMode::TOPK_NORMAL模式或者TopKMode::TOPK_NSMALL模式与kernel侧接口一致。isLargest输入表示降序/升序true表示降序false表示升序。与kernel侧接口一致。topKTiling输出输出TopK接口所需的tiling信息。返回值说明GetTopKMaxMinTmpSize返回值为true/falsetrue表示成功拿到TopK接口内部计算需要的最大和最小临时空间大小false表示获取失败。TopKTilingFunc返回值为true/falsetrue表示成功拿到TopK的Tiling各项参数值false表示获取失败。约束说明无调用示例如下样例介绍了使用TopK高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。将TopK Tiling结构体参数增加至TilingData结构体作为TilingData结构体的一个字段。namespace optiling { BEGIN_TILING_DATA_DEF(TilingData) TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, tilenum); // 添加其他tiling字段 ... TILING_DATA_FIELD_DEF(int32_t, k); TILING_DATA_FIELD_DEF(bool, islargest); TILING_DATA_FIELD_DEF(bool, isinitindex); TILING_DATA_FIELD_DEF(bool, ishasfinish); TILING_DATA_FIELD_DEF(uint32_t, tmpsize); TILING_DATA_FIELD_DEF(int32_t, outter); TILING_DATA_FIELD_DEF(int32_t, inner); TILING_DATA_FIELD_DEF(int32_t, n); TILING_DATA_FIELD_DEF(int32_t, order); TILING_DATA_FIELD_DEF(int32_t, sorted); TILING_DATA_FIELD_DEF_STRUCT(TopkTiling, topkTilingData); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(TopkCustom, TilingData) } // namespace optilingTiling实现函数中首先调用GetTopKMaxMinTmpSize接口获取TopK接口能完成计算所需最大/最小临时空间大小根据该范围结合实际的内存使用情况设置合适的空间大小然后根据输入shape等信息获取TopK kernel侧接口所需tiling参数。MERGE_SORT算法参考如下调用示例。namespace optiling { const uint32_t NUM_BLOCKS 8; const uint32_t TILE_NUM 8; const int32_t OUTTER 2; const int32_t INNER 32; const int32_t N 32; const int32_t K 8; const bool IS_LARGEST true; const bool IS_INITINDEX true; const bool IS_REUSESOURCE false; static ge::graphStatus TilingFunc(gert::TilingContext* context) { TilingData tiling; uint32_t totalLength context-GetInputTensor(0)-GetShapeSize(); context-SetBlockDim(NUM_BLOCKS); tiling.set_totalLength(totalLength); tiling.set_tileNum(TILE_NUM); tiling.set_k(K); tiling.set_outter(OUTTER); tiling.set_inner(INNER); tiling.set_n(N); tiling.set_islargest(IS_LARGEST); tiling.set_isinitindex(IS_INITINDEX); // 设置其他Tiling参数 ... // 本样例中仅作为样例说明通过GetTopKMaxMinTmpSize获取最小值并传入来保证功能正确开发者可以根据需要传入合适的空间大小。 uint32_t maxsize 0; uint32_t minsize 0; uint32_t dtypesize 4; // float类型 auto ascendcPlatform platform_ascendc::PlatformAscendC(context-GetPlatformInfo()); AscendC::TopKTilingFunc( ascendcPlatform, tiling.inner, tiling.outter, tiling.k, dtypesize, tiling.isinitindex, AscendC::TopKMode::TOPK_NSMALL, tiling.islargest, tiling.topkTilingData); AscendC::GetTopKMaxMinTmpSize( ascendcPlatform, tiling.inner, tiling.outter, IS_REUSESOURCE, tiling.isinitindex, AscendC::TopKMode::TOPK_NSMALL, tiling.islargest, dtypesize, maxsize, minsize); tiling.set_tmpsize(minsize); ... // 其他逻辑 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 optilingRADIX_SELECT算法参考如下调用示例。namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext* context) { std::mapge::DataType, uint32_t dtypeSizes {{ge::DataType::DT_UINT32, 4}, {ge::DataType::DT_INT32, 4}}; RadixtopkCustomTilingData tiling; const gert::RuntimeAttrs* attrs context-GetAttrs(); const uint32_t is_init_index *(attrs-GetAttrPointeruint32_t(0)); const uint32_t is_reuse_src *(attrs-GetAttrPointeruint32_t(1)); const uint32_t order *(attrs-GetAttrPointeruint32_t(2)); const uint32_t is_largest *(attrs-GetAttrPointeruint32_t(3)); const uint32_t outter *(attrs-GetAttrPointeruint32_t(4)); const uint32_t inner *(attrs-GetAttrPointeruint32_t(5)); const uint32_t n *(attrs-GetAttrPointeruint32_t(6)); const uint32_t k *(attrs-GetAttrPointeruint32_t(7)); const uint32_t k_pad *(attrs-GetAttrPointeruint32_t(8)); const uint32_t sorted *(attrs-GetAttrPointeruint32_t(9)); const uint32_t top_mode *(attrs-GetAttrPointeruint32_t(10)); auto xDType context-GetInputTensor(0)-GetDataType(); uint32_t typeSize dtypeSizes.at(xDType); AscendC::TopKConfig config; config.algo AscendC::TopKAlgo::RADIX_SELECT; if (order 1) { config.order AscendC::TopKOrder::LARGEST; } else if (order 2) { config.order AscendC::TopKOrder::SMALLEST; } else { config.order AscendC::TopKOrder::UNSET; } if (sorted 0) { config.sorted false; } else { config.sorted true; } uint32_t maxValue 0; uint32_t minValue 0; if (top_mode 0) { AscendC::GetTopKMaxMinTmpSize( inner, outter, k, is_reuse_src, is_init_index, AscendC::TopKMode::TOPK_NORMAL, is_largest, xDType, config, maxValue, minValue); context-SetTilingKey(0); } else { AscendC::GetTopKMaxMinTmpSize( inner, outter, k, is_reuse_src, is_init_index, AscendC::TopKMode::TOPK_NSMALL, is_largest, xDType, config, maxValue, minValue); context-SetTilingKey(1); } context-SetBlockDim(1); tiling.set_is_init_index(is_init_index); tiling.set_is_reuse_src(is_reuse_src); tiling.set_order(order); tiling.set_is_largest(is_largest); tiling.set_outter(outter); tiling.set_inner(inner); tiling.set_n(n); tiling.set_k(k); tiling.set_k_pad(k_pad); tiling.set_sorted(sorted); tiling.set_top_mode(top_mode); 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对应的kernel侧通过在核函数中调用GET_TILING_DATA获取TilingData继而将TilingData中的TopK Tiling信息传入TopK接口参与计算。完整的kernel侧样例请参考调用示例。extern C __global__ __aicore__ void topk_custom( GM_ADDR srcVal, GM_ADDR srcIdx, GM_ADDR finishLocal, GM_ADDR dstVal, GM_ADDR dstIdx, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelTopKfloat, true, true, false, false, AscendC::TopKMode::TOPK_NSMALL op; op.Init( srcVal, srcIdx, finishLocal, dstVal, dstIdx, tilingData.k, tilingData.islargest, tilingData.tmpsize, tilingData.outter, tilingData.inner, tilingData.n, tilingData.topkTilingData); op.Process(); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
相关文章
Linux Mint 22.3 Zena：轻量级LTS系统与老旧硬件优化指南
1. Linux Mint 22.3 Zena 版本概览Linux Mint 22.3 "Zena" 是2024年1月发布的长期支持版本（LTS），基于Ubuntu 22.04 Jammy Jellyfish构建。这个版本延续了Linux Mint一贯的"开箱即用"理念，预装了完整的媒体编解…
AscendC归约临时空间接口
GetSumMaxMinTmpSize 【免费下载链接】asc-devkit 本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。 项目地址: https://gi…
CANN Ascend C SIMD排序API文档
Sort 【免费下载链接】asc-devkit 本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。 项目地址: https://gitcode.com/cann/…
Ubuntu系统升级全攻略：从备份到优化
1. 升级前的准备工作在开始升级Ubuntu系统之前，有几个关键步骤需要完成。这些准备工作可以确保升级过程顺利，避免数据丢失或系统损坏。1.1 备份重要数据无论升级过程多么可靠，备份始终是必不可少的。我建议采用以下几种备份方式：文…
Tropos开源贡献指南：如何参与这个优秀天气项目的开发
Tropos开源贡献指南：如何参与这个优秀天气项目的开发 【免费下载链接】Tropos Weather and Forecasts for Humans 项目地址: https://gitcode.com/gh_mirrors/tr/Tropos Tropos是一个专注于为用户提供人性化天气和预报服务的开源项目。本文将为你提供一份完整…
Ethereum Security Toolbox部署指南：从本地开发到生产环境的完整配置
Ethereum Security Toolbox部署指南：从本地开发到生产环境的完整配置 【免费下载链接】eth-security-toolbox A Docker container preconfigured with all of the Trail of Bits Ethereum security tools. 项目地址: https://gitcode.com/gh_mirrors/et/eth-secur…
Flutter项目架构优化：使用json_model实现数据层的自动化管理
Flutter项目架构优化：使用json_model实现数据层的自动化管理 【免费下载链接】json_model Generate model class from Json file. 一行命令，通过Json文件生成Dart Model类。 项目地址: https://gitcode.com/gh_mirrors/js/json_model 在Flutter应…
GN+Ninja构建系统在RA4M2嵌入式开发中的应用
1. 为什么选择GNNinja构建系统在嵌入式开发领域，构建系统的选择直接影响开发效率和项目可维护性。传统Makefile虽然通用，但随着项目复杂度提升，其局限性日益明显。GN（Generate Ninja）作为元构建系统，配合Ni…
Ubuntu 24.04安装搜狗输入法的兼容性问题与解决方案
1. 为什么Ubuntu 24.04安装搜狗输入法需要特殊处理？ Ubuntu 24.04 LTS作为最新的长期支持版本，在输入法框架和显示服务器方面做出了重大调整。默认情况下，系统采用了Fcitx5输入法框架和Wayland显示协议，这与搜狗输入法基于Fcitx4的…
VS Code 高效配置与个性化定制全攻略
1. VS Code 高效配置基础作为一款轻量级但功能强大的代码编辑器，VS Code 的默认配置已经能满足基本需求，但通过合理调整设置可以大幅提升编码效率。我使用 VS Code 已经有五年多时间，期间尝试过各种配置方案，总结出这套适合大多数…
HarmonyOS 应用开发《掌上英语》第19篇:3D 翻转动画实现——ArkTS 动画系统全解析
3D 翻转动画实现——ArkTS 动画系统全解析引言 在移动应用中，卡片翻转动画是最受欢迎的交互动效之一，它能给用户带来直观的"物理世界"操作感。在我们的英语学习 App 的单词学习页面（CourseHomePage.ets）中，就…
BiSheng JDK-build性能调优：构建速度提升30%的优化策略
BiSheng JDK-build性能调优：构建速度提升30%的优化策略 【免费下载链接】bishengjdk-build BiSheng JDK build and test scripts - common across all releases/versions 项目地址: https://gitcode.com/openeuler/bishengjdk-build 前往项目官网免费下载&am…
MQ-2 烟雾传感器 ESP32-S3 驱动实战：ADC 读取与阈值校准 3 步完成
MQ-2 烟雾传感器 ESP32-S3 驱动实战：ADC 读取与阈值校准 3 步完成在智能家居和工业安全监测领域，烟雾检测是一个至关重要的环节。MQ-2 作为一款高性价比的半导体烟雾传感器，因其对多种可燃气体（如液化气、丙烷、氢气等）…
SPEC CPU 2006 v1.0.1 基准测试实战：ARM/X86/MIPS 三平台配置与 3 轮测试结果解读
SPEC CPU 2006 跨平台基准测试深度实战：ARM/X86/MIPS 架构配置优化与结果分析方法论在当今多元化的计算架构时代，如何客观评估不同处理器平台的真实性能成为系统工程师和性能优化专家的核心挑战。SPEC CPU 2006 作为业界公认的计算密集型基准测试套件&am…
每天60s读懂世界：2026年7月11日重点要闻解读
🔥 个人主页：杨利杰YJlio❄️ 个人专栏：《Windows 疑难杂症与工单复盘案例库》 《Sysinternals实战教程》《WINDOWS教程》 《Windows PowerShell 实战》 《人工智能实战合集》《超简单：用Python让Excel飞起来》🌟…
FAE放射组学分析工具：医学影像特征探索的完整解决方案
FAE放射组学分析工具：医学影像特征探索的完整解决方案 【免费下载链接】FAE FeAture Explorer 项目地址: https://gitcode.com/gh_mirrors/fae/FAE 你是否曾经面对海量医学影像数据感到无从下手？想要从CT、MRI等影像中提取有价值的定量特征&#…
0.69B参数实现中文多模态AI：揭秘Qwen3-SmVL模型融合技术的完整实战指南
0.69B参数实现中文多模态AI：揭秘Qwen3-SmVL模型融合技术的完整实战指南 【免费下载链接】happy-llm 📚 从零开始构建大模型 项目地址: https://gitcode.com/GitHub_Trending/ha/happy-llm 还在为大型多模态模型动辄数十亿参数、显存占用高而烦恼&…
解锁AMD Ryzen处理器性能潜力的SMU调试神器：从新手到专家的完整指南
解锁AMD Ryzen处理器性能潜力的SMU调试神器：从新手到专家的完整指南 【免费下载链接】SMUDebugTool A dedicated tool to help write/read various parameters of Ryzen-based systems, such as manual overclock, SMU, PCI, CPUID, MSR and Power Table. 项目地址…