游戏网站首页设计wordpress响应网页代码

游戏网站首页设计,wordpress响应网页代码,中国建设银行官方网站诚聘英才,网页设计图片地址怎么写基础知识准备 本文内容基于Ascend C算子开发衍生而来#xff0c;对于算子开发还不了解的读者可以通过以下资源进行学习#xff1a; 《Ascend C算子开发文档手册》#xff1a; https://www.hiascend.com/document/detail/zh/canncommercial/82RC1/opdevg/Ascendcopdevg/atl…基础知识准备本文内容基于Ascend C算子开发衍生而来对于算子开发还不了解的读者可以通过以下资源进行学习《Ascend C算子开发文档手册》https://www.hiascend.com/document/detail/zh/canncommercial/82RC1/opdevg/Ascendcopdevg/atlas_ascendc_10_0001.html《Ascend C算子开发系列课程入门》https://www.hiascend.com/developer/courses/detail/1691696509765107713《Ascend C算子开发系列课程进阶》https://www.hiascend.com/developer/courses/detail/1696414606799486977《Ascend C算子开发系列课程高级》https://www.hiascend.com/developer/courses/detail/1696690858236694530背景介绍随着越来越多的用户基于Ascend C进行算子开发算子开发过程中的易用性相关问题陆续被反馈出来其中多场景算子的高效开发实现以及维护管理一直是用户的痛点。本次在CANN全面开源开放中做了很多算子开发易用性提升的工作针对多场景算子的开发CANN推出了TilingKey模板化编程有效解决了多场景算子开发管理的问题。亮点介绍TilingKey模板化编程具备以下特点减少icache miss和scalar耗时提升算子调用性能通过模板化的开发方式提升算子开发效率算子代码的可读性和可维护性提升技术详细解析在算子开发过程中会遇到同一个算子有多种不同的实现的情形为了能够在运行时能够正确加载到对应的算子实现在Tiling的实现中会引入TilingKey的概念。TilingKey是一个算子内为了区分不同场景的实现而将Kernel代码进行区分的方法该方法类似于C的Template模板机制可减少不必要的icache miss以及scalar耗时有助于优化单次调用Kernel的性能。不同的Kernel实现分支可以通过TilingKey来标识Host侧设置TilingKey后可以选择对应的分支例如一个算子在不同的shape下有不同的算法逻辑Kernel侧可以通过TilingKey来选择不同的算法逻辑在Host侧Tiling算法也有差异Host/Kernel侧通过相同的TilingKey进行关联。假如有如下Kernel代码if (condition) { ProcessA(); } else { ProcessB(); }​如果函数ProcessA、ProcessB两个函数是个非常大的函数那么上述代码在编译后会变得更大而每次Kernel运行只会选择1个分支条件的判断和跳转在代码大到一定程度16-32K不同芯片存在差异后会出现icache miss。通过TilingKey可以对这种情况进行优化给2个Kernel的处理函数设置不同的TilingKey 1和2if (TILING_KEY_IS(1)) { ProcessA(); } elseif (TILING_KEY_IS(2)) { ProcessB(); }​这样Device Kernel编译时会自动识别到2个TilingKey并编译2个Kernel入口函数将条件判断进行常量折叠。同时需要和Host Tiling函数配合判断走ProcessA的场景设置TilingKey为1走ProcessB的场景设置TilingKey为2static ge::graphStatus TilingFunc(gert::TilingContext* context) { // some code if (condition) { context-SetTilingKey(1); } else { context-SetTilingKey(2); } return ge::GRAPH_SUCCESS; }​但是上述编程方式会让其他阅读这份代码的用户产生疑问假如用户从函数调用入口查看代码时看到了两个分支分别是TILING_KEY_IS(1)和TILING_KEY_IS(2)那么用户会无法理解“1”和“2”的含义。尤其是在涉及多个TilingKey的场景中开发者依赖TilingKey来管理Kernel的实现无论是在管理还是使用上都会遇到相当大的复杂性。因此为了简化这一过程消除TilingKey不易于记忆和理解的问题我们引入了一种采用模板编程的方法来替代传统的TilingKey编程的方案从而减少对TilingKey数值标识的依赖使Kernel的管理更加直观和高效。具体的开发步骤如下下面我们以add自定义算子开发为样例详细介绍一下具体的使用方法1.在自定义算子工程的op_kernel目录下新增定义模板参数和模板参数组合的头文件本示例中头文件命名为tiling_key_add_custom.h该头文件中需要包含模板头文件ascendc/host_api/tiling/template_argument.h。定义模板参数ASCENDC_TPL_ARGS_DECL和模板参数组合ASCENDC_TPL_ARGS_SEL即可使用的模板。具体API参考见模板参数定义。模板参数定义https://www.hiascend.com/document/detail/zh/canncommercial/82RC1/API/ascendcopapi/atlasascendc_api_07_00011.html#includeascendc/host_api/tiling/template_argument.h #define ADD_TPL_FP16 1 // 数据类型定义 #define ADD_TPL_FP32 0 #define ADD_TPL_ND 2 // 数据格式定义 #define ADD_TPL_NZ 29 // 模板参数 ASCENDC_TPL_ARGS_DECL(AddTemplateCustom, // 算子OpType ASCENDC_TPL_DTYPE_DECL(D_T_X, ADD_TPL_FP16, ADD_TPL_FP32), // DataType类型的模板参数定义输入参数x的数据类型取值范围为float16/float32 ASCENDC_TPL_DTYPE_DECL(D_T_Y, ADD_TPL_FP16, ADD_TPL_FP32), // DataType类型的模板参数定义输入参数y的数据类型取值范围为float16/float32 ASCENDC_TPL_DTYPE_DECL(D_T_Z, ADD_TPL_FP16, ADD_TPL_FP32), // DataType类型的模板参数定义输入参数z的数据类型取值范围为float16/float32 ASCENDC_TPL_UINT_DECL(TILE_NUM, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_MIX, 2, 0, 2, 3, 5, 10, 12, 13, 9, 8),// 自定义UINT类型无符号整形的模板参数定义模板参数为切分的块数编码位宽为ASCENDC_TPL_8_BW即8比特表示该模板参数的个数不超过8比特能表达的范围ASCENDC_TPL_UI_MIX表示通过混合模式表达取值范围有2组的数据{0-2}、{3-5}和穷举值10、12、13、9、8最后结果为{0, 1, 2, 3, 4, 5, 10, 12, 13, 9, 8} ASCENDC_TPL_BOOL_DECL(IS_SPLIT, 0, 1), // 自定义bool类型的模板参数定义模板参数为是否切分标志位取值范围为0和11表示切分0表示不切分 ); // 模板参数组合 // 用于调用GET_TPL_TILING_KEY获取TilingKey时接口内部校验TilingKey是否合法 ASCENDC_TPL_SEL( ASCENDC_TPL_ARGS_SEL( ASCENDC_TPL_DTYPE_SEL(D_T_X, ADD_TPL_FP16), ASCENDC_TPL_DTYPE_SEL(D_T_Y, ADD_TPL_FP16), ASCENDC_TPL_DTYPE_SEL(D_T_Z, ADD_TPL_FP16), ASCENDC_TPL_UINT_SEL(TILE_NUM, ASCENDC_TPL_UI_LIST, 1, 8), ASCENDC_TPL_BOOL_SEL(IS_SPLIT, 0, 1), ), ASCENDC_TPL_ARGS_SEL( ASCENDC_TPL_DTYPE_SEL(D_T_X, ADD_TPL_FP32), ASCENDC_TPL_DTYPE_SEL(D_T_Y, ADD_TPL_FP32), ASCENDC_TPL_DTYPE_SEL(D_T_Z, ADD_TPL_FP32), ASCENDC_TPL_UINT_SEL(TILE_NUM, ASCENDC_TPL_UI_LIST, 1, 8), ASCENDC_TPL_BOOL_SEL(IS_SPLIT, 0, 1), ), );2.Host侧调用GET_TPL_TILING_KEY接口生成TilingKeyGET_TPL_TILING_KEYhttps://www.hiascend.com/document/detail/zh/canncommercial/82RC1/API/ascendcopapi/atlas_ascendc_10_00008.htmlHost实现文件中包含步骤1中定义模板参数和模板参数组合的头文件。调用GET_TPL_TILING_KEY接口生成TilingKeyGET_TPL_TILING_KEY输入参数为模板参数的具体值传入时需要与定义模板参数和模板参数组合的头文件中的模板参数顺序保持一致。#include tiling_key_add_custom.h static ge::graphStatus TilingFunc(gert::TilingContext *context) { TilingData tiling; uint32_t totalLength context-GetInputShape(0)-GetOriginShape().GetShapeSize(); ge::DataType dtype_x context-GetInputDesc(0)-GetDataType(); ge::DataType dtype_y context-GetInputDesc(1)-GetDataType(); ge::DataType dtype_z context-GetOutputDesc(1)-GetDataType(); uint32_t D_T_X ADD_TPL_FP32, D_T_YADD_TPL_FP32, D_T_ZADD_TPL_FP32, TILE_NUM1, IS_SPLIT0; if(dtype_x ge::DataType::DT_FLOAT){ D_T_X ADD_TPL_FP32; }elseif(dtype_x ge::DataType::DT_FLOAT16){ D_T_X ADD_TPL_FP16; } if(dtype_y ge::DataType::DT_FLOAT){ D_T_Y ADD_TPL_FP32; }elseif(dtype_y ge::DataType::DT_FLOAT16){ D_T_Y ADD_TPL_FP16; } if(dtype_z ge::DataType::DT_FLOAT){ D_T_Z ADD_TPL_FP32; }elseif(dtype_z ge::DataType::DT_FLOAT16){ D_T_Z ADD_TPL_FP16; } if(totalLength MIN_LENGTH_FOR_SPLIT){ IS_SPLIT 0; TILE_NUM 1; }else{ IS_SPLIT 1; TILE_NUM DEFAULT_TILE_NUM; } context-SetBlockDim(BLOCK_DIM); tiling.set_totalLength(totalLength); tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity()); context-GetRawTilingData()-SetDataSize(tiling.GetDataSize()); constuint64_t tilingKey GET_TPL_TILING_KEY(D_T_X, D_T_Y, D_T_Z, TILE_NUM, IS_SPLIT); context-SetTilingKey(tilingKey); size_t *currentWorkspace context-GetWorkspaceSizes(1); currentWorkspace[0] 0; return ge::GRAPH_SUCCESS; }3.Kernel侧实现Kernel实现文件中包含步骤1中定义模板参数和模板参数组合的头文件。核函数添加template模板以便支持模板参数的传入参数顺序需要与定义模板参数和模板参数组合的头文件中的模板参数顺序保持一致。通过对模板参数的分支判断选择不同的Kernel侧实现。#include tiling_key_add_custom.h ... ... templateint D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); if(D_T_X ADD_TPL_FP32 D_T_Y ADD_TPL_FP32 D_T_Z ADD_TPL_FP32){ KernelAddfloat, float, float op; op.Init(x, y, z, tiling_data.totalLength, TILE_NUM, IS_SPLIT); op.Process1(); }else if(D_T_X ADD_TPL_FP16 D_T_Y ADD_TPL_FP16 D_T_Z ADD_TPL_FP16){ KernelAddhalf, half, half op; if(IS_SPLIT 0){ op.Init(x, y, z, tiling_data.totalLength, TILE_NUM, IS_SPLIT); op.Process1(); }else if(IS_SPLIT1){ op.Init(x, y, z, tiling_data.totalLength, TILE_NUM, IS_SPLIT); op.Process2(); } }}当前已开源的ops-nn仓和ops-transfomer仓中的部分核心算子已经基于该模板化编程方式完成优化感兴趣的同学可以访问ops-nn仓 BatMatMulV3算子和ops-transformer仓 FA算子。ops-nn仓 BatMatMulV3算子https://gitcode.com/cann/ops-nn/tree/master/matmul/batch_mat_mul_v3ops-transformer仓 FA算子https://gitcode.com/cann/ops-transformer/tree/master/attention/flash_attention_score总结TilingKey模板化编程既提供了一种简单的多场景算子开发编程范式简化了算子的开发难度也实际提升了算子的执行效率帮助开发者更便捷的开发高性能多场景算子。​​​​