已关闭
[Requirement|需求建议]: 新增 SIMD&SIMT 混合编程 UB 使用上限样例及测试覆盖 #1282
wulinyu创建于  7月18日关闭于  13 天前
wulinyu成员
7月18日 创建

Backgroud(背景信息)

新增 simd_simt_ub_usage_limit 样例工程,用于展示 SIMD&SIMT mixed 形态下 Unified Buffer(UB)的静态/动态内存使用方式、容量上限和越界问题定位方法。

该样例用于补齐现有 devkit 在 SIMD&SIMT 混合编程 UB 使用方式与容量边界方面的样例空缺。样例通过 Case 0~Case 3 讲述 UB 越界导致精度异常的定位方法、不同场景下的 UB 可用上限、UB 上的静态 UB 与动态 UB 布局,以及不同执行位置越界访问 UB 区域时的现象。

Origin(信息来源)

SIMT 需求访谈:补齐现有 devkit 样例空缺,方便开发者上手。

Benefit / Necessity (价值/作用)

帮助开发者理解混合编程场景下 SIMT VF、SIMD VF 与 MainScalar 对 UB 的共同使用关系,掌握静态内存、动态内存、预留空间和 Data Cache 对可用 UB 空间的影响。

Design(设计方案)

examples/05_simd_simt_hybrid/02_best_practices/simd_simt_ub_usage_limit/ 下新增样例工程,交付件包括:

  1. case0_ub_out_of_bounds.asc
  2. case1_ub_limit.asc
  3. case2_ub_layout.asc
  4. case3_out_of_bounds_result.asc
  5. CMakeLists.txt
  6. README.md / README_en.md

样例采用 SIMD&SIMT mixed 形态:__global__ __vector__ kernel 通过 asc_vf_call 调用 VF 函数。SIMT VF 使用 __simt_vf__,线程数配置为 1024;SIMD VF 使用 __simd_vf__。样例不引入 syncthreads() / syncwarp(),不专门验证 UB bank 冲突。

默认运行覆盖 4 个正向 场景:

  1. 无 SIMT VF 场景下,验证 SIMD/MainScalar 可使用的 UB 大小。
  2. 有 SIMT VF 场景下,验证用户可使用 UB 的大小上限,并体现 Data Cache 和预留空间影响。
  3. SIMT VF 与 SIMD VF 分别声明静态内存时,验证不同执行空间静态内存共同占用 UB 的容量限制。
  4. 使用 <<<...>>> 的动态 UB 参数配置动态内存,验证动态内存配置方式和访问范围。

样例通过 SCENARIO_NUM 覆盖 4 类 Case:

  1. Case 0:构造 UB 越界写入 Data Cache 导致输出精度异常的问题场景,展示执行状态确认、计算逻辑排查、访问范围计算和修复方法。
  2. Case 1:验证不调用 VF、仅调用 SIMD VF、仅调用 SIMT VF、同时调用 SIMD VF 和 SIMT VF 时的 UB 可用上限,并体现 Data Cache 和预留空间对上限的影响。
  3. Case 2:验证 MainScalar、SIMD VF 与 SIMT VF 分别声明静态 UB 时,不同执行位置的静态 UB 共同占用当前 Kernel 静态内存区域;同时说明动态 UB 共用 Kernel 启动时配置的动态 UB 空间。
  4. Case 3:构造 MainScalar 和 SIMT VF 越界访问 UB 区域的场景,展示越界访问未被占用的普通 UB 空间、预留空间和 Data Cache 时的运行结果和精度结果。

README 补充 UB 越界反例/调试 case:

  1. 说明静态内存、动态内存、预留空间和 Data Cache 的 UB 排布关系。
  2. 说明越界访问可能覆盖 Data Cache,导致结果不稳定或精度异常。
  3. 给出定位步骤:确认 Kernel 执行状态、核对 Host 侧 Golden 与 Kernel 计算逻辑、打印中间计算变量,并结合 UB 布局判断实际访问范围。
  4. 给出修复建议:保证 SIMT VF、SIMD VF 和 MainScalar 的静态/动态 UB 访问总范围不超过已申请空间,并为 Data Cache 保留空间。

STC 覆盖样例编译、默认运行、Case 0~Case 3 输出校验、Case 1 UB 边界、README 越界定位说明检视,以及默认路径不触发越界失败。

likedislike
Wwulinyu成员
7月18日 添加了label:requirement
Zzhangshunwei2成员
29 天前 关联了pull request:docs: 完善SIMT编程指导并新增UB使用限制样例
wulinyu成员
22 天前 评论:

需求开发中

likedislike
Zzhangshunwei2成员
15 天前 修改了issue 的描述
Zzhangshunwei2成员
15 天前 修改了issue 的描述
Zzhangshunwei2成员
15 天前 修改了issue 的描述
Zzhangshunwei2成员
15 天前 修改了issue 的描述
Zzhangshunwei2成员
13 天前 将 zhangshunwei2 设为负责人
Zzhangshunwei2成员
13 天前 issue状态由 待办的 改变为 已解决
Zzhangshunwei2成员
13 天前 关闭了 issue
CANN-robotCANN-robot成员
13 天前 添加了label:resolved