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:不支持
功能说明
本接口在SIMD和SIMT调试场景下提供assert断言功能。在算子Kernel侧的实现代码中,如果assert的内部条件判断不为真,则会输出assert条件,并将输入的信息格式化打印在屏幕上,同时算子运行失败。
在算子Kernel侧代码的适当位置使用assert进行断言检查,并格式化输出一些调试信息。示例如下:
C++
int assertFlag = 10;
assert(assertFlag != 10);
打印信息示例如下:
Text
[ASSERT] /home/.../add_custom.cpp:44: : Assertion `assertFlag != 10' failed.
</code></pre></div>
<p>请注意,assert接口的打印功能对算子的实际运行性能有影响。</p>
<h2>函数原型</h2>
<div class="code-block"><div class="code-header"><span class="lang-label">C++</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code><span class="n">assert</span><span class="p">(</span><span class="n">expr</span><span class="p">)</span>
</code></pre></div>
<h2>参数说明</h2>
<table>
<thead>
<tr>
<th>参数名</th>
<th>输入/输出</th>
<th>描述</th>
</tr>
</thead>
<tbody>
<tr>
<td>expr</td>
<td>输入</td>
<td>assert断言是否终止程序的条件。条件为true则程序继续执行,条件为false则终止程序。</td>
</tr>
</tbody>
</table>
<h2>返回值说明</h2>
<p>无</p>
<h2>约束说明</h2>
<ul>
<li>该接口当前仅支持融合编译场景。</li>
<li>SIMD场景,如果开发者需要包含标准库头文件<cassert>,请在"utils/debug/asc_assert.h"头文件之前包含,避免assert符号冲突。</li>
</ul>
<!-- npu="950,A3,910b" id13 -->
<ul>
<li>SIMT场景的支持情况如下:
<!-- npu="950" id10 -->
<ul>
<li>Ascend 950PR/Ascend 950DT:支持。
<!-- end id10 -->
<!-- npu="A3" id11 -->
</li>
<li>Atlas A3 训练系列产品/Atlas A3 推理系列产品:不支持。
<!-- end id11 -->
<!-- npu="910b" id12 -->
</li>
<li>Atlas A2 训练系列产品/Atlas A2 推理系列产品:不支持。
<!-- end id12 -->
</li>
</ul>
</li>
</ul>
<!-- end id13 -->
<h2>需要包含的头文件</h2>
<p>使用该接口需要包含"utils/debug/asc_assert.h"头文件。</p>
<div class="code-block"><div class="code-header"><span class="lang-label">C++</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code><span class="cp">#include</span><span class="w"> </span><span class="cpf">"utils/debug/asc_assert.h"</span>
</code></pre></div>
<h2>SIMD调用示例</h2>
<div class="code-block"><div class="code-header"><span class="lang-label">C++</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code><span class="c1">// SIMD</span>
<span class="n">__global__</span><span class="w"> </span><span class="n">__cube__</span><span class="w"> </span><span class="kt">void</span><span class="w"> </span><span class="n">simp_test_equal</span><span class="p">(</span><span class="kt">int</span><span class="w"> </span><span class="n">a</span><span class="p">)</span>
<span class="p">{</span>
<span class="w"> </span><span class="c1">// input a is 7</span>
<span class="w"> </span><span class="n">assert</span><span class="p">(</span><span class="n">a</span><span class="w"> </span><span class="o">==</span><span class="w"> </span><span class="mi">6</span><span class="p">);</span>
<span class="p">}</span>
</code></pre></div>
<p>程序运行时会触发assert,打印效果如下:</p>
<div class="code-block"><div class="code-header"><span class="lang-label">Text</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code>[ASSERT] test_assert.asc:41: : Assertion `a == 6' failed.
SIMT调用示例
SIMT编程场景:
C++__global__ __launch_bounds__(1024) inline void simt_kernel(float* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; assert(!isnan(x[idx])); }程序运行时会触发assert,打印效果如下:
Text[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(float *): Assertion `!isnan(x[idx])' failed. </code></pre></div> </li> <li> <p>SIMD与SIMT混合编程场景:</p> <div class="code-block"><div class="code-header"><span class="lang-label">C++</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code><span class="n">__simt_vf__</span><span class="w"> </span><span class="nf">__launch_bounds__</span><span class="p">(</span><span class="mi">1024</span><span class="p">)</span><span class="w"> </span><span class="kr">inline</span><span class="w"> </span><span class="kt">void</span><span class="w"> </span><span class="n">simt_kernel</span><span class="p">(</span><span class="n">__gm__</span><span class="w"> </span><span class="kt">float</span><span class="o">*</span><span class="w"> </span><span class="n">x</span><span class="p">)</span> <span class="p">{</span> <span class="w"> </span><span class="kt">int</span><span class="w"> </span><span class="n">idx</span><span class="w"> </span><span class="o">=</span><span class="w"> </span><span class="n">threadIdx</span><span class="p">.</span><span class="n">x</span><span class="w"> </span><span class="o">+</span><span class="w"> </span><span class="n">blockIdx</span><span class="p">.</span><span class="n">x</span><span class="w"> </span><span class="o">*</span><span class="w"> </span><span class="n">blockDim</span><span class="p">.</span><span class="n">x</span><span class="p">;</span> <span class="w"> </span><span class="n">assert</span><span class="p">(</span><span class="o">!</span><span class="n">isnan</span><span class="p">(</span><span class="n">x</span><span class="p">[</span><span class="n">idx</span><span class="p">]));</span> <span class="p">}</span> </code></pre></div> <p>程序运行时会触发assert,打印效果如下:</p> <div class="code-block"><div class="code-header"><span class="lang-label">Text</span><button class="copy-btn" title="复制代码"><svg width="16" height="16" viewBox="0 0 16 16"><path d="M0 6.75C0 5.784.784 5 1.75 5h1.5a.75.75 0 0 1 0 1.5h-1.5a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-1.5a.75.75 0 0 1 1.5 0v1.5A1.75 1.75 0 0 1 9.25 16h-7.5A1.75 1.75 0 0 1 0 14.25Z"/><path d="M5 1.75C5 .784 5.784 0 6.75 0h7.5C15.216 0 16 .784 16 1.75v7.5A1.75 1.75 0 0 1 14.25 11h-7.5A1.75 1.75 0 0 1 5 9.25Zm1.75-.25a.25.25 0 0 0-.25.25v7.5c0 .138.112.25.25.25h7.5a.25.25 0 0 0 .25-.25v-7.5a.25.25 0 0 0-.25-.25Z"/></svg></button></div><pre class="highlight"><code>[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(__gm__ float *): Assertion `!isnan(x[idx])' failed.