12a14b50创建于 9月15日历史提交

assert/ascendc_assert

产品支持情况

Ascend 950PR/Ascend 950DT：支持

Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持

Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持

Atlas 200I/500 A2 推理产品：不支持

Atlas 推理系列产品AI Core：支持

Atlas 推理系列产品Vector Core：支持

Atlas 训练系列产品：不支持

Kirin X90：不支持

Kirin 9030：不支持

不同产品对各编程场景的支持情况如下表所示。

表1 支持能力

产品	SIMD	SIMT VF	SIMD VF
Ascend 950PR/Ascend 950DT	支持	支持	支持
Atlas A3 训练系列产品/Atlas A3 推理系列产品	支持	不支持	不支持
Atlas A2 训练系列产品/Atlas A2 推理系列产品	支持	不支持	不支持
Atlas 推理系列产品 AI Core	支持	不支持	不支持
Atlas 推理系列产品 Vector Core	支持	不支持	不支持

功能说明

头文件路径为："utils/debug/asc_assert.h"。

assert和ascendc_assert用于在Device侧代码中进行断言检查。断言条件成立时，后续代码继续执行；断言条件不成立时，接口打印断言失败信息并触发异常，算子执行失败。

需要在断言失败时输出自定义错误信息，可以在条件后传入格式字符串和对应参数。assert和ascendc_assert提供相同的断言能力，可以任选其一使用。

注意

该接口用于调测，使用时会影响算子性能，生产环境建议通过设置ASCENDC_DUMP=0关闭。

函数原型

assert(expr)
assert(expr, fmt, args...)

ascendc_assert(expr)
ascendc_assert(expr, fmt, args...)

参数说明

表2 参数说明

参数名	输入/输出	描述
expr	输入	断言条件。为`true`时，核函数（Kernel）继续运行；为`false`时，接口打印断言失败信息并触发异常。
fmt	输入	可选。自定义错误信息的格式控制字符串，包含普通字符和转换说明。普通字符原样输出；转换说明以百分号（%）开始，用于控制`args`中对应参数的转换和输出。支持的转换类型和数据类型参见printf的参数说明。
args	输入	可选。与`fmt`中的转换说明对应的附加参数。每个参数替换`fmt`中对应的转换说明，参数数量和类型应与转换说明匹配。

返回值说明

无

约束说明

通用约束

在CMake工程的Release模式下，默认使能NDEBUG宏定义，本接口不生效。
使用assert时，如需包含标准库头文件<cassert>，请先包含<cassert>，再包含"kernel_operator.h"或"utils/debug/asc_assert.h"，避免标准库覆盖Ascend C的assert定义。

SIMD编程场景

CPU域调试时，多参数调用仅检查断言条件，不打印fmt和args指定的自定义错误信息，行为与单参数调用相同。
单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。超出限制时，断言失败信息不会打印，但接口仍会触发异常。可以通过aclInit接口中的simd_printf_fifo_size_per_core字段配置，配置范围为[1KB,64MB]。PyTorch调用和算子入图场景暂不支持该配置。

SIMT VF编程场景

不支持CPU域调试。
不输出fmt和args指定的自定义错误信息。

SIMD VF编程场景

CPU域调试时，多参数调用仅检查断言条件，不打印fmt和args指定的自定义错误信息，行为与单参数调用相同。
fmt可以直接使用字符串字面量。使用字符串指针时，该指针必须指向Unified Buffer（UB）中的字符串，示例如下：
```
__ubuf__ const char* fmt = "value is %d.\n";
assert(value == 6, fmt, value);
```
每个AIV核在单次asc_vf_call执行期间使用2KB预留UB空间临时保存调测数据。同一次asc_vf_call中的assert、ascendc_assert、printf和asc_dump共享该空间。该空间中的数据传输完成后会被复用，因此上述接口产生的累计调测数据可以超过2KB。单条断言失败信息必须能完整保存在该空间中，否则该条信息不会打印，但接口仍会触发异常。
SIMD VF场景下，本接口需要使用2KB预留UB空间传递断言信息。开启--cce-disable-asc-reserved-ubuf选项后，本接口不可用。
SIMD VF场景下，simd_printf_fifo_size_per_core建议配置为3KB以上。配置过小且打印数据量较大时，断言失败信息不会被打印，但接口仍会触发异常。

调用示例

SIMD编程场景

条件不满足时，接口打印AIC/AIV类型、Block号、源码位置、函数名和条件表达式。传入格式字符串和对应参数时，还会打印自定义错误信息。随后，接口触发异常。

以下示例中，value为7，不满足value == 6，因此触发断言。

extern "C" __global__ __vector__ void AssertSimdKernel()
{
    int32_t value = 7;
    AscendC::printf("before assert, value is %d.\n", value);
    assert(value == 6, "value is %d.\n", value);
}

程序运行时会触发断言，实际打印如下：

[AIV Block 0/1] before assert, value is 7.
[AIV Block 0/1] [ASSERT] .../assert_simd.asc:10: void AssertSimdKernel(): Assertion 'value == 6' failed. value is 7.

SIMT VF编程场景

条件不满足时，接口打印源码位置、函数名和条件表达式，然后触发异常。

SIMT编程场景：

__global__ __launch_bounds__(1024) inline void simt_kernel(float* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    assert(!isnan(x[idx]));
}

程序运行时会触发assert，打印效果如下：

[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(float *): Assertion `!isnan(x[idx])' failed.

SIMD与SIMT混合编程场景：

__simt_vf__ __launch_bounds__(1024) inline void simt_kernel(__gm__ float* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    assert(!isnan(x[idx]));
}

程序运行时会触发assert，打印效果如下：

[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(__gm__ float *): Assertion `!isnan(x[idx])' failed.

SIMD VF编程场景

条件不成立时，接口打印源码位置、函数名和条件表达式；如果传入了格式字符串和对应参数，自定义错误信息会追加在该断言失败信息之后。随后，接口触发异常。

__simd_vf__ inline void AssertSimdVf()
{
    int32_t value = 7;
    assert(value == 6, "value is %d.\n", value);
}

程序运行时会触发断言，打印格式如下：

[ASSERT] .../assert_simd_vf.asc:12: void AssertSimdVf(): Assertion 'value == 6' failed. value is 7.