Merged
feat: add asc_uintdiv SIMT API #6009
zhangshunwei2Add on Sep 14
feat: add asc_uintdiv SIMT API #6009
Merged
Merged fromzhangshunwei2/asc-devkitfastdiv_simtintocann/asc-devkitmaster
Total: 7 files changed+302-1
| @@ -3113,6 +3113,8 @@ | |||
| 3113 | - [\_\_umulhi](SIMT-API/math_functions/integer_math_functions/__umulhi.md) | 3113 | - [\_\_umulhi](SIMT-API/math_functions/integer_math_functions/__umulhi.md) |
| 3114 | - [\_\_mul64hi](SIMT-API/math_functions/integer_math_functions/__mul64hi.md) | 3114 | - [\_\_mul64hi](SIMT-API/math_functions/integer_math_functions/__mul64hi.md) |
| 3115 | - [\_\_umul64hi](SIMT-API/math_functions/integer_math_functions/__umul64hi.md) | 3115 | - [\_\_umul64hi](SIMT-API/math_functions/integer_math_functions/__umul64hi.md) |
| 3116 | + - [asc_uintdiv](SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) | ||
| 3117 | + - [asc_get_uintdiv_magic_and_shift](SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) | ||
| 3116 | - [\_\_mul\_i32toi64](SIMT-API/math_functions/integer_math_functions/__mul_i32toi64.md) | 3118 | - [\_\_mul\_i32toi64](SIMT-API/math_functions/integer_math_functions/__mul_i32toi64.md) |
| 3117 | - [\_\_brev](SIMT-API/math_functions/integer_math_functions/__brev.md) | 3119 | - [\_\_brev](SIMT-API/math_functions/integer_math_functions/__brev.md) |
| 3118 | - [\_\_clz](SIMT-API/math_functions/integer_math_functions/__clz.md) | 3120 | - [\_\_clz](SIMT-API/math_functions/integer_math_functions/__clz.md) |
| @@ -783,6 +783,8 @@ | |||
| 783 | | [__umulhi](../math_functions/integer_math_functions/__umulhi.md) | 获取输入uint32类型数据x和y乘积的高32位。 | | 783 | | [__umulhi](../math_functions/integer_math_functions/__umulhi.md) | 获取输入uint32类型数据x和y乘积的高32位。 | |
| 784 | | [__mul64hi](../math_functions/integer_math_functions/__mul64hi.md) | 获取输入int64类型数据x和y乘积的高64位。 | | 784 | | [__mul64hi](../math_functions/integer_math_functions/__mul64hi.md) | 获取输入int64类型数据x和y乘积的高64位。 | |
| 785 | | [__umul64hi](../math_functions/integer_math_functions/__umul64hi.md) | 获取输入uint64类型数据x和y乘积的高64位。 | | 785 | | [__umul64hi](../math_functions/integer_math_functions/__umul64hi.md) | 获取输入uint64类型数据x和y乘积的高64位。 | |
| 786 | +| [asc_uintdiv](../math_functions/integer_math_functions/asc_uintdiv.md) | 使用magic和shift参数执行快速无符号除法。 | | ||
| 787 | +| [asc_get_uintdiv_magic_and_shift](../math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) | 根据无符号除数计算快速除法所需的magic和shift参数。 | | ||
| 786 | | [__mul_i32toi64](../math_functions/integer_math_functions/__mul_i32toi64.md) | 计算输入32位整数x和y的乘积，返回64位结果。 | | 788 | | [__mul_i32toi64](../math_functions/integer_math_functions/__mul_i32toi64.md) | 计算输入32位整数x和y的乘积，返回64位结果。 | |
| 787 | | [__brev](../math_functions/integer_math_functions/__brev.md) | 将输入数据的位序反转，返回反转后的值。 | | 789 | | [__brev](../math_functions/integer_math_functions/__brev.md) | 将输入数据的位序反转，返回反转后的值。 | |
| 788 | | [__clz](../math_functions/integer_math_functions/__clz.md) | 从输入数据的二进制最高有效位开始，返回连续的前导零的位数。 | | 790 | | [__clz](../math_functions/integer_math_functions/__clz.md) | 从输入数据的二进制最高有效位开始，返回连续的前导零的位数。 | |
Adocs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md+116-0
| @@ -0,0 +1,116 @@ | |||
| 1 | +# asc_get_uintdiv_magic_and_shift | ||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +<!-- npu="950" id1 --> | ||
| 6 | +- Ascend 950PR/Ascend 950DT：支持 | ||
| 7 | +<!-- end id1 --> | ||
| 8 | +<!-- npu="A3" id2 --> | ||
| 9 | +- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持 | ||
| 10 | +<!-- end id2 --> | ||
| 11 | +<!-- npu="910b" id3 --> | ||
| 12 | +- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持 | ||
| 13 | +<!-- end id3 --> | ||
| 14 | +<!-- npu="310b" id4 --> | ||
| 15 | +- Atlas 200I/500 A2 推理产品：不支持 | ||
| 16 | +<!-- end id4 --> | ||
| 17 | +<!-- npu="310p" id5 --> | ||
| 18 | +- Atlas 推理系列产品AI Core：不支持 | ||
| 19 | +- Atlas 推理系列产品Vector Core：不支持 | ||
| 20 | +<!-- end id5 --> | ||
| 21 | +<!-- npu="910" id6 --> | ||
| 22 | +- Atlas 训练系列产品：不支持 | ||
| 23 | +<!-- end id6 --> | ||
| 24 | +<!-- @ref: asc-devkit/res/docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift_res.md#id1 --> | ||
| 25 | + | ||
| 26 | +## 功能说明 | ||
| 27 | + | ||
| 28 | +头文件路径为：`"simt_api/math_functions.h"`。 | ||
| 29 | + | ||
| 30 | +整数快除用于同除数场景的快速计算，算法分为两步计算商，本接口用于第一步计算，获取乘法因子`magic`和移位量`shift`，生成的参数供SIMT VF侧快除接口[asc_uintdiv](asc_uintdiv.md)使用。本接口的计算过程如下： | ||
| 31 | + | ||
| 32 | +$$ | ||
| 33 | +\begin{aligned} | ||
| 34 | +shift &= \left\lceil \log_2 divisor \right\rceil \\ | ||
| 35 | +magic &= \left\lfloor 2^N \times \frac{2^{shift} - divisor}{divisor} \right\rfloor + 1 | ||
| 36 | +\end{aligned} | ||
| 37 | +$$ | ||
| 38 | + | ||
| 39 | +根据除数数据类型的不同，公式中$N$的取值不同：当除数为`uint32_t`数据类型时，$N$为32；当除数为`uint64_t`数据类型时，$N$为64。 | ||
| 40 | + | ||
| 41 | +整数快除算法可以用于性能优化，适用于一个除数被用于多次除法计算的场景。 | ||
| 42 | + | ||
| 43 | +> [!NOTE]说明 | ||
| 44 | +> | ||
| 45 | +> 整数快除算法使用方式： | ||
| 46 | +> 1. 在标量计算逻辑中，调用`asc_get_uintdiv_magic_and_shift`接口，计算除数的`magic`和`shift`值。 | ||
| 47 | +> 2. 在SIMT VF计算逻辑中，调用[asc_uintdiv](asc_uintdiv.md)接口，完成指定被除数的除法运算。 | ||
| 48 | + | ||
| 49 | +> [!CAUTION]注意 | ||
| 50 | +> | ||
| 51 | +> 本接口不能在SIMT VF中调用，只能在标量计算逻辑中调用。 | ||
| 52 | + | ||
| 53 | +## 函数原型 | ||
| 54 | + | ||
| 55 | +- 除数为`uint32_t`类型 | ||
| 56 | + | ||
| 57 | + ```cpp | ||
| 58 | + inline void asc_get_uintdiv_magic_and_shift(uint32_t* magic, uint32_t* shift, uint32_t divisor) | ||
| 59 | + ``` | ||
| 60 | + | ||
| 61 | +- 除数为`uint64_t`类型 | ||
| 62 | + | ||
| 63 | + ```cpp | ||
| 64 | + inline void asc_get_uintdiv_magic_and_shift(uint64_t* magic, uint64_t* shift, uint64_t divisor) | ||
| 65 | + ``` | ||
| 66 | + | ||
| 67 | +## 参数说明 | ||
| 68 | + | ||
| 69 | +**表1** 参数说明 | ||
| 70 | + | ||
| 71 | +| 参数名 | 输入/输出 | 描述 | | ||
| 72 | +| --- | --- | --- | | ||
| 73 | +| magic | 输出 | 根据指定除数计算得到的快速除法中的乘法因子。 | | ||
| 74 | +| shift | 输出 | 根据指定除数计算得到的快速除法中的移位量。 | | ||
| 75 | +| divisor | 输入 | 除数。 | | ||
| 76 | + | ||
| 77 | +## 返回值说明 | ||
| 78 | + | ||
| 79 | +无 | ||
| 80 | + | ||
| 81 | +## 约束说明 | ||
| 82 | + | ||
| 83 | +- 当除数为`uint32_t`数据类型时，支持的除数最大值为`INT32_MAX`，即2^31-1；当除数为`uint64_t`数据类型时，支持的除数最大值为`INT64_MAX`，即2^63-1。 | ||
| 84 | +- 除数不支持为0，否则计算结果未定义。 | ||
| 85 | +- SIMT编程场景当前不支持使用该接口。 | ||
| 86 | + | ||
| 87 | +## 调用示例 | ||
| 88 | + | ||
| 89 | +完整样例请参考[SIMD与SIMT混合编程实现快速除法算子样例](../../../../../../examples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README.md)。 | ||
| 90 | + | ||
| 91 | +- SIMD与SIMT混合编程场景： | ||
| 92 | + | ||
| 93 | + ```cpp | ||
| 94 | + #include "simt_api/asc_simt.h" | ||
| 95 | + #include "simt_api/math_functions.h" | ||
| 96 | + | ||
| 97 | + __simt_vf__ __launch_bounds__(1024) inline void simt_fast_div( | ||
| 98 | + __gm__ uint32_t* output, uint32_t value, uint32_t magic, uint32_t shift, uint32_t total_length) | ||
| 99 | + { | ||
| 100 | + uint32_t idx = threadIdx.x; | ||
| 101 | + if (idx >= total_length) { | ||
| 102 | + return; | ||
| 103 | + } | ||
| 104 | + output[idx] = asc_uintdiv(value, magic, shift); | ||
| 105 | + } | ||
| 106 | + | ||
| 107 | + __global__ __vector__ void integer_div_kernel( | ||
| 108 | + __gm__ uint32_t* output, uint32_t value, uint32_t divisor, uint32_t total_length) | ||
| 109 | + { | ||
| 110 | + // 在标量计算逻辑中，根据除数预计算magic和shift。 | ||
| 111 | + uint32_t magic = 0; | ||
| 112 | + uint32_t shift = 0; | ||
| 113 | + asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor); | ||
| 114 | + asc_vf_call<simt_fast_div>(dim3(1024), output, value, magic, shift, total_length); | ||
| 115 | + } | ||
| 116 | + ``` | ||
| @@ -0,0 +1,109 @@ | |||
| 1 | +# asc_uintdiv | ||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +<!-- npu="950" id1 --> | ||
| 6 | +- Ascend 950PR/Ascend 950DT：支持 | ||
| 7 | +<!-- end id1 --> | ||
| 8 | +<!-- npu="A3" id2 --> | ||
| 9 | +- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持 | ||
| 10 | +<!-- end id2 --> | ||
| 11 | +<!-- npu="910b" id3 --> | ||
| 12 | +- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持 | ||
| 13 | +<!-- end id3 --> | ||
| 14 | +<!-- npu="310b" id4 --> | ||
| 15 | +- Atlas 200I/500 A2 推理产品：不支持 | ||
| 16 | +<!-- end id4 --> | ||
| 17 | +<!-- npu="310p" id5 --> | ||
| 18 | +- Atlas 推理系列产品AI Core：不支持 | ||
| 19 | +- Atlas 推理系列产品Vector Core：不支持 | ||
| 20 | +<!-- end id5 --> | ||
| 21 | +<!-- npu="910" id6 --> | ||
| 22 | +- Atlas 训练系列产品：不支持 | ||
| 23 | +<!-- end id6 --> | ||
A | |||
| 24 | +<!-- @ref: asc-devkit/res/docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv_res.md#id1 --> | ||
| 25 | + | ||
| 26 | +## 功能说明 | ||
| 27 | + | ||
| 28 | +头文件路径为：`"simt_api/math_functions.h"`。 | ||
| 29 | + | ||
| 30 | +整数快除算法分为两步计算商`quotient`。第一步根据固定除数预计算`magic`和`shift`，可由配套预计算接口[asc_get_uintdiv_magic_and_shift](asc_get_uintdiv_magic_and_shift.md)完成；第二步调用本接口，根据被除数`dividend`、`magic`和`shift`计算商。计算公式如下： | ||
| 31 | + | ||
| 32 | +$$ | ||
| 33 | +quotient = \left(\left\lfloor\frac{dividend \times magic}{2^N}\right\rfloor + dividend\right) \gg shift | ||
| 34 | +$$ | ||
| 35 | + | ||
| 36 | +根据被除数数据类型的不同，公式中$N$的取值不同：当被除数为`uint32_t`数据类型时，$N$为32；当被除数为`uint64_t`数据类型时，$N$为64。 | ||
| 37 | + | ||
| 38 | +整数快除算法可以用于性能优化，适用于一个除数被用于多次除法计算的场景。 | ||
| 39 | + | ||
| 40 | +> [!NOTE]说明 | ||
| 41 | +> | ||
| 42 | +> 整数快除算法使用方式： | ||
| 43 | +> 1. 在标量计算逻辑中，调用[asc_get_uintdiv_magic_and_shift](asc_get_uintdiv_magic_and_shift.md)接口，计算除数的`magic`和`shift`值。 | ||
| 44 | +> 2. 在SIMT VF计算逻辑中，调用`asc_uintdiv`接口，完成指定被除数的除法运算。 | ||
| 45 | + | ||
| 46 | +## 函数原型 | ||
| 47 | + | ||
| 48 | +- 被除数为`uint32_t`类型 | ||
| 49 | + | ||
| 50 | + ```cpp | ||
| 51 | + inline uint32_t asc_uintdiv(uint32_t dividend, uint32_t magic, uint32_t shift) | ||
| 52 | + ``` | ||
| 53 | + | ||
| 54 | +- 被除数为`uint64_t`类型 | ||
| 55 | + | ||
| 56 | + ```cpp | ||
| 57 | + inline uint64_t asc_uintdiv(uint64_t dividend, uint64_t magic, uint64_t shift) | ||
| 58 | + ``` | ||
| 59 | + | ||
| 60 | +## 参数说明 | ||
| 61 | + | ||
| 62 | +**表1** 参数说明 | ||
| 63 | + | ||
| 64 | +| 参数名 | 输入/输出 | 描述 | | ||
| 65 | +| --- | --- | --- | | ||
| 66 | +| dividend | 输入 | 被除数。 | | ||
| 67 | +| magic | 输入 | 根据指定除数预计算得到的快速除法中的乘法因子。 | | ||
| 68 | +| shift | 输入 | 根据指定除数预计算得到的快速除法中的移位量。 | | ||
| 69 | + | ||
| 70 | +## 返回值说明 | ||
| 71 | + | ||
| 72 | +返回`dividend`除以`magic`和`shift`所对应的除数得到的商。 | ||
| 73 | + | ||
| 74 | +## 约束说明 | ||
| 75 | + | ||
| 76 | +- 当被除数为`uint32_t`数据类型时，支持的被除数最大值为`INT32_MAX`，即2^31-1；当被除数为`uint64_t`数据类型时，支持的被除数最大值为`INT64_MAX`，即2^63-1。 | ||
| 77 | +- `magic`和`shift`必须对应同一非零除数；`uint32_t`数据和`uint64_t`数据对应除数的最大值分别为`INT32_MAX`和`INT64_MAX`。 | ||
| 78 | +- SIMT编程场景当前不支持使用该接口。 | ||
| 79 | + | ||
| 80 | +## 调用示例 | ||
| 81 | + | ||
| 82 | +完整样例请参考[SIMD与SIMT混合编程实现快速除法算子样例](../../../../../../examples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README.md)。 | ||
| 83 | + | ||
| 84 | +- SIMD与SIMT混合编程场景： | ||
| 85 | + | ||
| 86 | + ```cpp | ||
| 87 | + #include "simt_api/asc_simt.h" | ||
| 88 | + #include "simt_api/math_functions.h" | ||
| 89 | + | ||
| 90 | + __simt_vf__ __launch_bounds__(1024) inline void simt_fast_div( | ||
| 91 | + __gm__ uint32_t* output, uint32_t value, uint32_t magic, uint32_t shift, uint32_t total_length) | ||
| 92 | + { | ||
| 93 | + uint32_t idx = threadIdx.x; | ||
| 94 | + if (idx >= total_length) { | ||
| 95 | + return; | ||
| 96 | + } | ||
| 97 | + output[idx] = asc_uintdiv(value, magic, shift); | ||
| 98 | + } | ||
| 99 | + | ||
| 100 | + __global__ __vector__ void integer_div_kernel( | ||
| 101 | + __gm__ uint32_t* output, uint32_t value, uint32_t divisor, uint32_t total_length) | ||
| 102 | + { | ||
| 103 | + // 在标量计算逻辑中，根据除数预计算magic和shift。 | ||
| 104 | + uint32_t magic = 0; | ||
| 105 | + uint32_t shift = 0; | ||
| 106 | + asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor); | ||
| 107 | + asc_vf_call<simt_fast_div>(dim3(1024), output, value, magic, shift, total_length); | ||
| 108 | + } | ||
| 109 | + ``` | ||
| @@ -24,6 +24,10 @@ | |||
| 24 | 24 | ||
| 25 | - **[\_\_umul64hi](__umul64hi.md)** | 25 | - **[\_\_umul64hi](__umul64hi.md)** |
| 26 | 26 | ||
| 27 | +- **[asc_uintdiv](asc_uintdiv.md)** | ||
| 28 | + | ||
| 29 | +- **[asc_get_uintdiv_magic_and_shift](asc_get_uintdiv_magic_and_shift.md)** | ||
| 30 | + | ||
| 27 | - **[\_\_mul\_i32toi64](__mul_i32toi64.md)** | 31 | - **[\_\_mul\_i32toi64](__mul_i32toi64.md)** |
| 28 | 32 | ||
| 29 | - **[\_\_brev](__brev.md)** | 33 | - **[\_\_brev](__brev.md)** |
| @@ -55,4 +59,3 @@ | |||
| 55 | - **[max](max.md)** | 59 | - **[max](max.md)** |
| 56 | 60 | ||
| 57 | - **[min](min.md)** | 61 | - **[min](min.md)** |
| 58 | - | ||
| @@ -5932,6 +5932,65 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline float norm4df(float a, float b, float c, f | |||
| 5932 | 5932 | ||
| 5933 | #endif // ASCENDC_USE_LEGACY_PRECISION | 5933 | #endif // ASCENDC_USE_LEGACY_PRECISION |
| 5934 | 5934 | ||
| 5935 | +#ifndef __NPU_COMPILER_INTERNAL_PURE_SIMT__ | ||
| 5936 | +#ifndef ASCENDC_CPU_DEBUG | ||
| 5937 | +__SIMT_DEVICE_FUNCTIONS_DECL__ inline uint32_t asc_uintdiv(uint32_t dividend, uint32_t magic, uint32_t shift) | ||
| 5938 | +{ | ||
| 5939 | + uint32_t q = __umulhi(dividend, magic); | ||
| 5940 | + uint32_t sum = dividend + q; | ||
| 5941 | + return sum >> shift; | ||
| 5942 | +} | ||
| 5943 | + | ||
| 5944 | +__SIMT_DEVICE_FUNCTIONS_DECL__ inline uint64_t asc_uintdiv(uint64_t dividend, uint64_t magic, uint64_t shift) | ||
| 5945 | +{ | ||
| 5946 | + uint64_t q = __umul64hi(dividend, magic); | ||
| 5947 | + uint64_t sum = dividend + q; | ||
| 5948 | + return sum >> shift; | ||
| 5949 | +} | ||
| 5950 | + | ||
| 5951 | +__aicore__ inline uint64_t asc_get_uintdiv_magic_impl(uint64_t dividend, uint64_t divisor) | ||
| 5952 | +{ | ||
| 5953 | + uint64_t quotient = 0; | ||
| 5954 | + uint64_t remainder = dividend; | ||
| 5955 | + for (int i = 0; i < 64; ++i) { // 64: process every bit of the uint64_t dividend. | ||
| 5956 | + quotient <<= 1; | ||
| 5957 | + const uint64_t borrow = remainder >> 63; // 63: highest bit index of uint64_t. | ||
| 5958 | + remainder <<= 1; | ||
| 5959 | + if (borrow != 0) { | ||
| 5960 | + remainder -= divisor; | ||
| 5961 | + quotient |= 1; // 1: set the current quotient bit. | ||
| 5962 | + } else if (remainder >= divisor) { | ||
| 5963 | + remainder -= divisor; | ||
| 5964 | + quotient |= 1; // 1: set the current quotient bit. | ||
| 5965 | + } | ||
| 5966 | + } | ||
| 5967 | + return quotient + 1; // 1: round the reciprocal up for exact division. | ||
| 5968 | +} | ||
| 5969 | + | ||
| 5970 | +__aicore__ inline void asc_get_uintdiv_magic_and_shift(uint32_t* magic, uint32_t* shift, uint32_t divisor) | ||
| 5971 | +{ | ||
| 5972 | + const int64_t pos = 64 - clz(static_cast<uint64_t>(divisor)); // 64: width of the clz input. | ||
| 5973 | + const int64_t cnt1 = bcnt1(static_cast<uint64_t>(divisor)); | ||
| 5974 | + *shift = cnt1 == 1 ? pos - 1 : pos; | ||
| 5975 | + *magic = (1ULL << 32) * ((1ULL << *shift) - divisor) / divisor + 1; // 32: uint32_t bit width. | ||
| 5976 | +} | ||
| 5977 | + | ||
| 5978 | +__aicore__ inline void asc_get_uintdiv_magic_and_shift(uint64_t* magic, uint64_t* shift, uint64_t divisor) | ||
| 5979 | +{ | ||
| 5980 | + const int64_t pos = 64 - clz(divisor); // 64: uint64_t bit width. | ||
| 5981 | + const int64_t cnt1 = bcnt1(divisor); | ||
| 5982 | + *shift = cnt1 == 1 ? pos - 1 : pos; | ||
| 5983 | + uint64_t dividend = 0; | ||
| 5984 | + if (*shift < 64) { // 64: one-past-the-highest valid uint64_t shift. | ||
| 5985 | + dividend = (1ULL << *shift) - divisor; | ||
| 5986 | + } else { | ||
| 5987 | + dividend = UINT64_MAX - divisor + 1; // 1: two's-complement form of 2^64 - divisor. | ||
| 5988 | + } | ||
| 5989 | + *magic = asc_get_uintdiv_magic_impl(dividend, divisor); | ||
| 5990 | +} | ||
| 5991 | +#endif | ||
| 5992 | +#endif | ||
| 5993 | + | ||
| 5935 | #endif | 5994 | #endif |
| 5936 | #endif // IMPL_SIMT_API_MATH_FUNCTIONS_IMPL_H | 5995 | #endif // IMPL_SIMT_API_MATH_FUNCTIONS_IMPL_H |
| 5937 | 5996 | ||
| @@ -259,6 +259,16 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline float fdividef(float x, float y); | |||
| 259 | 259 | ||
| 260 | __SIMT_DEVICE_FUNCTIONS_DECL__ inline int signbit(float x); | 260 | __SIMT_DEVICE_FUNCTIONS_DECL__ inline int signbit(float x); |
| 261 | 261 | ||
| 262 | +#ifndef __NPU_COMPILER_INTERNAL_PURE_SIMT__ | ||
| 263 | +__SIMT_DEVICE_FUNCTIONS_DECL__ inline uint32_t asc_uintdiv(uint32_t dividend, uint32_t magic, uint32_t shift); | ||
| 264 | + | ||
| 265 | +__SIMT_DEVICE_FUNCTIONS_DECL__ inline uint64_t asc_uintdiv(uint64_t dividend, uint64_t magic, uint64_t shift); | ||
| 266 | + | ||
| 267 | +__aicore__ inline void asc_get_uintdiv_magic_and_shift(uint32_t* magic, uint32_t* shift, uint32_t divisor); | ||
| 268 | + | ||
| 269 | +__aicore__ inline void asc_get_uintdiv_magic_and_shift(uint64_t* magic, uint64_t* shift, uint64_t divisor); | ||
| 270 | +#endif | ||
| 271 | + | ||
| 262 | #ifndef __NPU_COMPILER_INTERNAL_PURE_SIMT__ | 272 | #ifndef __NPU_COMPILER_INTERNAL_PURE_SIMT__ |
| 263 | #if (__NPU_ARCH__ == 3510) | 273 | #if (__NPU_ARCH__ == 3510) |
| 264 | #ifndef ASCENDC_CPU_DEBUG | 274 | #ifndef ASCENDC_CPU_DEBUG |
[建议] [文档规范] docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md:23 同目录 28 个 API 页均在 npu 标签块之后带有 资源引用注释行，本文档缺失（若该行由资源流水线自动生成，可忽略本条）。 建议：补充 @ref 注释行，与同目录其他 API 文档保持一致。