SinCosCompute性能调优样例【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit概述本样例以sincos计算为例介绍Ascend C SIMT编程方式下的线程配置优化思路。样例包含1个基线版本以及1个优化版本基线版本中未设置__launch_bounds__编译器按照默认值1024即每个Block内1024个线程进行资源分配导致寄存器溢出优化版本通过配置__launch_bounds__(512)提示编译器每个Block的最大线程数量为512编译器根据提示在编译过程中充分利用硬件资源从而避免寄存器溢出展示SIMT编程方式下合理配置线程数优化性能的调优路径。支持的产品Ascend 950PR/Ascend 950DT支持的CANN软件版本 CANN 9.1.0目录结构介绍sincos_compute/ │ ├── CMakeLists.txt // cmake编译文件 │ ├── sincos_compute.asc // sincos样例实现 │ └── README.md样例描述样例功能使用sincosf函数同时计算sin和cos结果sincosf(input[idx], output_sin idx, output_cos idx)样例规格样例类型(OpType)SinCosCompute样例输入nameshapedata typeformatinput[393216]floatND样例输出output_sin[393216]floatNDoutput_cos[393216]floatND核函数名sincos_thread_1024 / sincos_thread_512样例实现sincos计算说明本样例使用Ascend C提供的sincosf接口同时计算sin/cos结果样例中设置了固定shape每个线程计算16个输入值计算流程如下根据算子数据shape的切分逻辑计算每个核的起始地址调用Ascend C提供的sincosf接口同时计算sin/cos结果线程数与寄存器关系在SIMT编程模型中核函数定义时配置的最大线程数直接影响每个线程可用的寄存器数量最大线程数每个线程可用寄存器个数1025~204816513~102432257~512641~256127关键原则配置的最大线程数越大每个线程可用寄存器数越少计算密集型算子单个线程所需的寄存器通常较多一般建议配置512或1024线程数据搬运类算子单个线程所需的寄存器通常较少一般建议配置2048线程当寄存器不足以存下所有的临时变量时会出现寄存器溢出register spill数据会溢出到栈空间Global Memory导致性能下降样例实现说明本样例通过2个独立的kernel来体现__launch_bounds__的效果每个kernel对应特定的Case版本。Case实现特点使用的核函数优化特性Case 0不设置launch bounds使用默认值sincos_thread_1024基线版本未配置__launch_bounds__Case 1根据实际算子的计算规模配置__launch_bounds__(512)sincos_thread_512编译器使用用户指定的配置值进行相应优化性能指标说明指标说明Task Duration(us)整个任务执行的总时间算子执行时间以该参数为准DCache Read GMDCache从Global Memory读取数据的次数DCache Read VectorVector Core从DCache读取数据的次数DCache Write VectorVector Core向DCache写入数据的次数Case 0: 基线版本寄存器溢出样例目标不配置__launch_bounds__观察寄存器溢出对性能的影响核心实现默认线程数为1024每个Thread仅可用32个寄存器sincosf计算需要更多寄存器超出32个寄存器限制关键代码__global__ void sincos_thread_1024(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }编译信息使用--cce-res-usage编译选项查看寄存器使用情况[BISHENG] Function properties for _Z18sincos_thread_1024PfS_S_m_simt_entry: Stack size: 32 bytes, Used register number: 32分析Stack size: 32 bytes → 存在寄存器溢出Used register number: 32 → 达到1024线程下的寄存器上限寄存器溢出导致中间数据存储到Global Memory增加访存开销性能数据Task Duration(us)DCache Read GMDCache Read VectorDCache Write Vector102.47256640768性能瓶颈寄存器溢出到Global Memory额外的栈空间访问增加延迟DCache Read Vector / DCache Write Vector次数较高640次 / 768次优化方向通过__launch_bounds__提示编译器真实的blockDim充分利用寄存器资源避免寄存器溢出。Case 1: 优化版本避免寄存器溢出优化目标通过配置__launch_bounds__(512)避免寄存器溢出充分利用寄存器资源提升性能核心优化指定__launch_bounds__(512)每个Thread可用64个寄存器sincosf计算所需的寄存器在限制范围内关键代码__global__ __launch_bounds__(512) void sincos_thread_512(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }编译信息[BISHENG] Function properties for _Z17sincos_thread_512PfS_S_m_simt_entry: Stack size: 0 bytes, Used register number: 48分析Stack size: 0 bytes → 无寄存器溢出Used register number: 48 → 在64个寄存器限制内所有中间数据保存在寄存器避免Global Memory访问性能数据Task Duration(us)DCache Read GM(次)DCache Read Vector(次)DCache Write Vector(次)96.22256512256优化效果分析Task Duration从102.47us降低到96.22us耗时下降约6.1%DCache Read GM保持不变说明并没有增加额外开销DCache Read Vector从640减小至512 DCache Write Vector从768减小至256说明没有寄存器溢出后对于Data Cache的读写次数减少stack物理位置位于Global Memory因此寄存器溢出时对于stack的访问会体现在Data Cache的访问次数上寄存器充分利用避免数据溢出到Global Memory性能对比总结Ascend 950PR性能数据综合优化效果从Case 0基线版本到Case 1优化版本Task Duration从102.47us降低到96.22us耗时下降约6.1%DCache Read Vector从640减小至512DCache Write Vector从768减小至256Case versionTask Duration(us)Task Duration相对Case 0优化点Case 0102.471x基线版本寄存器溢出到Global MemoryCase 196.220.94x耗时配置launch bounds避免寄存器溢出调优建议识别寄存器溢出使用--cce-res-usage编译选项查看寄存器使用情况Stack size 0存在寄存器溢出Stack size 0无寄存器溢出合理配置线程数计算密集型算子建议512或1024线程数据搬运类算子建议2048线程根据寄存器需求表选择合适的线程数配置使用__launch_bounds__提示编译器__global__ __launch_bounds__(线程数) void kernel_name(...)验证优化效果对比优化前后的编译信息Stack size和Used register number对比优化前后的性能数据Task Duration、DCache Read GM、DCache Read Vector、DCache Write Vector编译运行在本样例根目录下执行如下步骤编译并执行样例。配置环境变量请根据当前环境上CANN开发套件包的安装方式配置环境变量。source ${install_path}/cann/set_env.sh说明${install_path}为CANN包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。样例执行mkdir -p build cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-3510 ..;make -j; # 编译工程 ./sincos_compute 1024 # 执行基线样例 ./sincos_compute 512 # 执行优化样例编译选项说明选项可选值说明CMAKE_ASC_ARCHITECTURESdav-3510NPU 架构本样例仅支持 dav-3510Ascend 950PR/Ascend 950DT执行结果如下说明精度对比成功。[Success] Case accuracy is verification passed.性能分析使用msprof工具获取详细性能数据msprof op ./sincos_compute 1024 # 分析基线case的性能 msprof op ./sincos_compute 512 # 分析优化case的性能命令完成后会在默认目录下生成以OPPROF_{timestamp}_XXX命名的文件夹,性能数据文件夹结构示例如下├──dump # 原始的性能数据用户无需关注 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比 ├──L2Cache.csv # L2 Cache命中率 ├──Memory.csv # UBL1和主存储器读写带宽速率 ├──MemoryL0.csv # L0AL0B和L0C读写带宽速率 ├──MemoryUB.csv # Vector和Scalar到UB的读写带宽速率 ├──OpBasicInfo.csv # 算子基础信息 ├──PipeUtilization.csv # 采集计算单元和搬运单元耗时和占比 ├──ResourceConflictRatio.csv # UB上的 bank group、bank conflict和资源冲突率在所有指令中的占比 └──visualize_data.bin # MindStudio Insight呈现文件查看具体的性能分析结果# 如查看算子基础信息 cat ./OPPROF_*/OpBasicInfo.csv # 如查看内存相关数据 cat ./OPPROF_*/Memory.csv【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考