Background（背景信息）

• predicate 优化是控制流优化的重要手段，可避免同一 warp 内出现不同的执行路径，从而在控制流指令（if、switch、do、for、while）引发分支发散时，消除不同分支串行执行带来的分支切换开销。
• 本需求旨在补齐 SIMT predicate 优化样例，覆盖两类典型分支场景，对同一分支链路采用不同处理方式对比性能差异，展示 predicate 优化的收益来源与适用边界，帮助用户理解其底层逻辑，在算子开发中对分支结构做出合理取舍。

Origin（信息来源）

SIMT 样例开发需求

Benefit / Necessity （价值/作用）

• 帮助用户理解 SIMT 编程模式下 Warp Divergence 的性能影响。
• 展示 predicate 优化的收益来源与适用边界。
• 说明不同分支命中率场景下应如何组合使用提前 return 与 predicate 优化。
• 提供可编译、可运行、可测量的性能对比样例，便于用户复现和参考。

Design（设计方案）

编程模型

• 线程索引：threadIdx.x + blockIdx.x * blockDim.x。
• 执行配置：blocks_per_grid = 8，THREAD_COUNT = 1024，共 8192 线程。
  入参与出参
• 入参：输入数组，case0和case1为 half 类型，case2、case3 和 case4 为 float 类型，长度均为 8192。每组 case 用同一套输入分批运行，按输入分布分为以下数据批：
  • hlog10 组：全正常值批（不含特殊值）、全特殊值批（六个特殊值均匀填充）、混合批。
  • j0f 组：低值段为主批（ax <= 8.0f）、高值段为主批（ax > 8.0f，含中值段）、全特殊值批（NaN / ±Inf）、混合批。
• 出参：输出数组，长度 8192；hlog10 组为 half 类型，j0f 组为 float 类型。
  case对照设计.
  case
  分支结构
  对照范围
  case0（hlog10）

特殊值处理放在正常值路径之前，命中即提前 return
基线

case1（hlog10）

特殊值处理放在正常值路径之后，统一 return 临时变量 temp
优化方式，与基线对照特殊值处理的位置

case2（j0f）
NaN/Inf 特殊值处理放在正常数据之前，所有分支命中即直接返回
基线

case3（j0f）

正常数据分段处理后逐级覆盖临时变量 result，NaN/Inf 特殊值处理放在最后，统一 return
优化方式一：全部分支统一 return，与 case2 对照返回方式
case4（j0f）
覆盖范围广的低值段与中值段命中后直接返回，高值段算出 result，特殊值处理放在最后覆盖 result，统一 return
优化方式二：大范围分支直接返回、剩余统一 return，与 case2 对照返回方式

计算逻辑
case0（hlog10）：特殊值处理放在正常值路径之前，命中即 return
先对精度有偏差的特殊值做矫正，命中即返回修正结果；全部未命中则继续执行正常值路径。
SIMT_DEVICE_FUNCTIONS_DECL inline half hlog10_case0(half x)
{
if (x == static_cast(0.2362060546875f)) {
return static_cast(-0.62646484375f);
} else if (x == static_cast(0.2490234375f)) {
return static_cast(-0.60400390625f);
} else if (x == static_cast(3.0703125f)) {
return static_cast(0.487060546875f);
} else if (x == static_cast(126.0625f)) {
return static_cast(2.099609375f);
} else if (x == static_cast(11496.0f)) {
return static_cast(4.05859375f);
} else if (x == static_cast(2976.0f)) {
return static_cast(3.474609375f);
}
float x_fp32 = __cvt_float<__internal_get_round<__RoundMode::CAST_RINT>(), RoundingSaturation::RS_DISABLE_VALUE>(x);
float result_fp32 = logf(x_fp32) / logf(10.0f); // log10(x) = log(x) / log(10)
return __cvt_half<__internal_get_round<__RoundMode::CAST_RINT>(), RoundingSaturation::RS_DISABLE_VALUE>(result_fp32);
}
case1（hlog10）：特殊值处理放在正常值路径之后，统一 return temp
先完整执行正常值路径得到 temp，再对精度有偏差的特殊值做矫正，覆盖 temp，最后统一返回 temp。
SIMT_DEVICE_FUNCTIONS_DECL inline half hlog10_case1(half x)
{
float x_fp32 = __cvt_float<__internal_get_round<__RoundMode::CAST_RINT>(), RoundingSaturation::RS_DISABLE_VALUE>(x);
float result_fp32 = logf(x_fp32) / logf(10.0f); // log10(x) = log(x) / log(10)
half temp = __cvt_half<__internal_get_round<__RoundMode::CAST_RINT>(), RoundingSaturation::RS_DISABLE_VALUE>(result_fp32);
if (x == static_cast(0.2362060546875f)) {
temp = static_cast(-0.62646484375f);
} else if (x == static_cast(0.2490234375f)) {
temp = static_cast(-0.60400390625f);
} else if (x == static_cast(3.0703125f)) {
temp = static_cast(0.487060546875f);
} else if (x == static_cast(126.0625f)) {
temp = static_cast(2.099609375f);
} else if (x == static_cast(11496.0f)) {
temp = static_cast(4.05859375f);
} else if (x == static_cast(2976.0f)) {
temp = static_cast(3.474609375f);
}
return temp;
}
case2（j0f）：所有分支命中即直接返回
先处理 NaN/Inf 特殊值，命中即返回；再依次处理低值段、中值段、高值段，任一段命中即返回该段结果，无统一出口。
SIMT_DEVICE_FUNCTIONS_DECL inline float j0f_case2(float x)
{
if (__isnan(x)) {
return x;
}
float ax = __fabsf(x);
if (__isinf(ax)) {
return 0.0f;
}
if (ax <= 8.0f) {
return __internal_j0f_less8(ax);
}
if (ax <= 1.0e13f) {
return __internal_j0f_middle_range(ax);
}
return __internal_j0f_huge_range(ax);
}
case3（j0f）：全部分支命中后逐级覆盖，统一 return
先按数值最大的高值段算出 result，随后中值段与低值段处理依次执行，命中则覆盖 result，最后经 NaN/Inf 特殊值处理覆盖 result，统一返回。
SIMT_DEVICE_FUNCTIONS_DECL inline float j0f_case3(float x)
{
float ax = __fabsf(x);
float result = __internal_j0f_huge_range(ax);
if (ax <= 1.0e13f) {
result = __internal_j0f_middle_range(ax);
}
if (ax <= 8.0f) {
result = __internal_j0f_less8(ax);
}
if (__isnan(x)) {
result = x; // NaN -> NaN
}
if (__isinf(ax)) {
result = 0.0f; // ±Inf -> 0
}
return result;
}
case4（j0f）：部分分支直接返回，剩余统一 return
先处理覆盖范围广的低值段与中值段，命中即返回；未命中的输入按高值段算出 result，再经 NaN/Inf 特殊值处理覆盖 result，最后统一返回。
SIMT_DEVICE_FUNCTIONS_DECL inline float j0f_case4(float x)
{
float ax = __fabsf(x);
if (ax <= 8.0f) {
return __internal_j0f_less8(ax);
}
if (ax <= 1.0e13f) {
return __internal_j0f_middle_range(ax);
}
float result = __internal_j0f_huge_range(ax);
if (__isnan(x)) {
result = x;
}
if (__isinf(ax)) {
result = 0.0f;
}
return result;
}
预期结果

• hlog10 组：
  • 全正常值批：case1 略快于 case0。
  • 全特殊值批：case0 快于 case1。
  • 混合批：两者接近，差异由特殊值占比决定。
• j0f 组：
  • 低值段为主批（ax <= 8.0f）：case4 最快，case2 次之，case3 最慢。
  • 高值段为主批（ax > 8.0f）：case2 与 case4 接近且都快于 case3。
  • 全特殊值批：case2 最快，case3 次之，case4 最慢。
  • 混合批：case4 总体最优。
    出口标准
• 版本：主线。
• 芯片：Ascend 950，架构 dav-3510。
• 软件环境：CANN ≥ 9.1.0。
• 交付件：predicate_optimization.asc、CMakeLists.txt、README.md、SRS、STC、 branch_handing.md。
• 功能：case0 与 case1 对比 hlog10 特殊值处理两种写法的性能，case2、case3 与 case4 对比 j0f 分支结构三种写法的性能，按输入分布分批统计正确性、cycle 与 Task Duration。
• 工程：样例工程可按 README 完成配置和编译，Host 侧逐数据批调用各 case kernel 并汇总打印正确性与耗时结果。
• 测试：STC 包含 5 个 case 的测试用例，初始执行状态为未执行；当前执行结果待补充。