news 2026/6/10 20:33:24

cann/asc-devkit SinCosCompute性能调优样例

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
cann/asc-devkit SinCosCompute性能调优样例

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 typeformat
    input[393216]floatND
    样例输出output_sin[393216]floatND
    output_cos[393216]floatND
    核函数名sincos_thread_1024 / sincos_thread_512

样例实现

sincos计算说明

本样例使用Ascend C提供的sincosf接口同时计算sin/cos结果，样例中设置了固定shape，每个线程计算16个输入值，计算流程如下：

  1. 根据算子数据shape的切分逻辑，计算每个核的起始地址
  2. 调用Ascend C提供的sincosf接口，同时计算sin/cos结果

线程数与寄存器关系

在SIMT编程模型中，核函数定义时配置的最大线程数直接影响每个线程可用的寄存器数量：

最大线程数每个线程可用寄存器个数
1025~204816
513~102432
257~51264
1~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 Vector
102.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减小至512，DCache Write Vector从768减小至256
Case versionTask Duration(us)Task Duration相对Case 0优化点
Case 0102.471x基线版本，寄存器溢出到Global Memory
Case 196.220.94x耗时配置launch bounds避免寄存器溢出

调优建议

  1. 识别寄存器溢出：使用--cce-res-usage编译选项查看寄存器使用情况

    • Stack size > 0：存在寄存器溢出
    • Stack size = 0：无寄存器溢出

  2. 合理配置线程数：

    • 计算密集型算子：建议512或1024线程
    • 数据搬运类算子：建议2048线程
    • 根据寄存器需求表选择合适的线程数配置

  3. 使用__launch_bounds__提示编译器：

    __global__ __launch_bounds__(线程数) void kernel_name(...)

  4. 验证优化效果：

    • 对比优化前后的编译信息（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_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程 ./sincos_compute 1024 # 执行基线样例 ./sincos_compute 512 # 执行优化样例

  • 编译选项说明

    选项可选值说明
    CMAKE_ASC_ARCHITECTURESdav-3510NPU 架构：本样例仅支持 dav-3510（Ascend 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 # UB，L1和主存储器读写带宽速率 ├──MemoryL0.csv # L0A，L0B，和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），仅供参考

版权声明: 本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若内容造成侵权/违法违规/事实不符，请联系邮箱：809451989@qq.com进行投诉反馈，一经查实，立即删除！
网站建设 2026/6/10 20:29:26

探索scodec核心组件：BitVector与Codec trait深度剖析 [特殊字符]

探索scodec核心组件&#xff1a;BitVector与Codec trait深度剖析 &#x1f680; 【免费下载链接】scodec Scala combinator library for working with binary data 项目地址: https://gitcode.com/gh_mirrors/sc/scodec 在Scala生态系统中处理二进制数据时&#xff0c;s…

作者头像 李华
网站建设 2026/6/10 20:28:02

Heroku Buildpack for Elixir缓存机制揭秘：加速依赖安装与应用构建

Heroku Buildpack for Elixir缓存机制揭秘&#xff1a;加速依赖安装与应用构建 【免费下载链接】heroku-buildpack-elixir Heroku Buildpack for Elixir with nitro boost 项目地址: https://gitcode.com/gh_mirrors/he/heroku-buildpack-elixir 你是否曾经为Elixir应用…

作者头像 李华
网站建设 2026/6/10 20:23:41

Code2Prompt：AI代码分析与智能重构的上下文工程解决方案

Code2Prompt&#xff1a;AI代码分析与智能重构的上下文工程解决方案 【免费下载链接】code2prompt A CLI tool to convert your codebase into a single LLM prompt with source tree, prompt templating, and token counting. 项目地址: https://gitcode.com/GitHub_Trendin…

作者头像 李华
网站建设 2026/6/10 20:23:22

Logseq移动端深度评测：如何在手机上构建你的第二大脑？

Logseq移动端深度评测&#xff1a;如何在手机上构建你的第二大脑&#xff1f; 【免费下载链接】logseq A privacy-first, open-source platform for knowledge management and collaboration. Download link: http://github.com/logseq/logseq/releases. roadmap: https://logs…

作者头像 李华
网站建设 2026/6/10 20:20:54

Rusty V8完整指南：5步掌握在Rust中运行JavaScript

Rusty V8完整指南&#xff1a;5步掌握在Rust中运行JavaScript 【免费下载链接】rusty_v8 Rust bindings for the V8 JavaScript engine 项目地址: https://gitcode.com/gh_mirrors/ru/rusty_v8 Rusty V8是一个将强大的V8 JavaScript引擎无缝集成到Rust应用中的高质量绑定…

作者头像 李华