Star399
273
代码
代码
Issues12
Pull Requests51
Actions
项目讨论
Wiki
分析
Star399
273
已合并
[需求]TDIV TDIVS TRECIP add high precision algorithm #648
小呀小石头创建于 3月31日
[需求]TDIV TDIVS TRECIP add high precision algorithm #648
已合并
小呀小石头创建于 3月31日
从denglei0610/pto-isamaster合入到cann/pto-isamaster
共 28 个文件变更+347-266
Mdocs/isa/TDIV.md+9-1
@@ -41,10 +41,16 @@ pto.tdiv ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst :
41Declared in `include/pto/common/pto_instr.hpp`:41Declared in `include/pto/common/pto_instr.hpp`:
42 42 
43```cpp43```cpp
44-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>44+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
45+ typename TileDataSrc1, typename... WaitEvents>
45PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events);46PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events);
46```47```
47 48 
49+`PrecisionType` has the following values available:
50+ 
51+* `DivAlgorithm::DEFAULT`: Normal algorithm, faster but with lower precision.
52+* `DivAlgorithm::HIGH_PRECISION`: High precision algorithm, but slower.
53+ 
48## Constraints54## Constraints
49 55 
50- **Implementation checks (A2A3)**:56- **Implementation checks (A2A3)**:
@@ -63,6 +69,8 @@ PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &sr
63 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain;.69 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain;.
64- **Division-by-zero**:70- **Division-by-zero**:
65 - Behavior is target-defined.71 - Behavior is target-defined.
72+- **High Precision Algorithm**
73+ - Only available on A5, `PrecisionType` option is ignored on A3.
66 74 
67## Examples75## Examples
68 76 
Mdocs/isa/TDIVS.md+17-4
@@ -55,13 +55,22 @@ pto.tdivs ins(%scalar, %src : dtype, !pto.tile_buf<...>) outs(%dst : !pto.tile_b
55Declared in `include/pto/common/pto_instr.hpp`:55Declared in `include/pto/common/pto_instr.hpp`:
56 56 
57```cpp57```cpp
58-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>58+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
59-PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar, WaitEvents &... events);59+ typename... WaitEvents>
60+PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar,
61+ WaitEvents &... events);
60 62 
61-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>63+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
62-PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0, WaitEvents &... events);64+ typename... WaitEvents>
65+PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0,
66+ WaitEvents &... events)
63```67```
64 68 
69+`PrecisionType` has the following values available:
70+ 
71+* `DivAlgorithm::DEFAULT`: Normal algorithm, faster but with lower precision.
72+* `DivAlgorithm::HIGH_PRECISION`: High precision algorithm, but slower.
73+ 
65## Constraints74## Constraints
66 75 
67- **Implementation checks (A2A3)** (both overloads):76- **Implementation checks (A2A3)** (both overloads):
@@ -80,6 +89,8 @@ PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar,
80 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain.89 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain.
81- **Division-by-zero**:90- **Division-by-zero**:
82 - Behavior is target-defined; on A5 the tile/scalar form maps to multiply-by-reciprocal and uses `1/0 -> +inf` for `scalar == 0`.91 - Behavior is target-defined; on A5 the tile/scalar form maps to multiply-by-reciprocal and uses `1/0 -> +inf` for `scalar == 0`.
92+- **High Precision Algorithm**
93+ - Only available on A5, `PrecisionType` option is ignored on A3.
83 94 
84## Examples95## Examples
85 96 
@@ -94,6 +105,7 @@ void example_auto() {
94 using TileT = Tile<TileType::Vec, float, 16, 16>;105 using TileT = Tile<TileType::Vec, float, 16, 16>;
95 TileT src, dst;106 TileT src, dst;
96 TDIVS(dst, src, 2.0f);107 TDIVS(dst, src, 2.0f);
108+ TDIVS<DivAlgorithm::HIGH_PRECISION>(dst, src, 2.0f);
97}109}
98```110```
99 111 
@@ -110,6 +122,7 @@ void example_manual() {
110 TASSIGN(src, 0x1000);122 TASSIGN(src, 0x1000);
111 TASSIGN(dst, 0x2000);123 TASSIGN(dst, 0x2000);
112 TDIVS(dst, 2.0f, src);124 TDIVS(dst, 2.0f, src);
125+ TDIVS<DivAlgorithm::HIGH_PRECISION>(dst, 2.0f, src);
113}126}
114```127```
115 128 
Mdocs/isa/TDIVS_zh.md+17-4
@@ -55,13 +55,22 @@ pto.tdivs ins(%scalar, %src : dtype, !pto.tile_buf<...>) outs(%dst : !pto.tile_b
55声明于 `include/pto/common/pto_instr.hpp`：55声明于 `include/pto/common/pto_instr.hpp`：
56 56 
57```cpp57```cpp
58-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>58+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
59-PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar, WaitEvents &... events);59+ typename... WaitEvents>
60+PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar,
61+ WaitEvents &... events);
60 62 
61-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>63+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
62-PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0, WaitEvents &... events);64+ typename... WaitEvents>
65+PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0,
66+ WaitEvents &... events)
63```67```
64 68 
69+`PrecisionType`可指定以下值：
70+ 
71+* `DivAlgorithm::DEFAULT`：普通算法，速度快但精度较低。
72+* `DivAlgorithm::HIGH_PRECISION`：高精度算法，速度较慢。
73+ 
65## 约束74## 约束
66 75 
67- **实现检查 (A2A3)**（两个重载）:76- **实现检查 (A2A3)**（两个重载）:
@@ -85,6 +94,8 @@ PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar,
85 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域.94 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域.
86- **除零**:95- **除零**:
87 - 行为由目标定义；在 A5 上，tile/标量形式映射到乘以倒数，并对 `scalar == 0` 使用 `1/0 -> +inf`。96 - 行为由目标定义；在 A5 上，tile/标量形式映射到乘以倒数，并对 `scalar == 0` 使用 `1/0 -> +inf`。
97+- **高精度算法**
98+ - 仅在A5上有效，`PrecisionType`选项A3上将被忽略。
88 99 
89## 示例100## 示例
90 101 
@@ -99,6 +110,7 @@ void example_auto() {
99 using TileT = Tile<TileType::Vec, float, 16, 16>;110 using TileT = Tile<TileType::Vec, float, 16, 16>;
100 TileT src, dst;111 TileT src, dst;
101 TDIVS(dst, src, 2.0f);112 TDIVS(dst, src, 2.0f);
113+ TDIVS<DivAlgorithm::HIGH_PRECISION>(dst, src, 2.0f);
102}114}
103```115```
104 116 
@@ -115,6 +127,7 @@ void example_manual() {
115 TASSIGN(src, 0x1000);127 TASSIGN(src, 0x1000);
116 TASSIGN(dst, 0x2000);128 TASSIGN(dst, 0x2000);
117 TDIVS(dst, 2.0f, src);129 TDIVS(dst, 2.0f, src);
130+ TDIVS<DivAlgorithm::HIGH_PRECISION>(dst, 2.0f, src);
118}131}
119```132```
120 133 
Mdocs/isa/TDIV_zh.md+11-1
@@ -41,10 +41,16 @@ pto.tdiv ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst :
41声明于 `include/pto/common/pto_instr.hpp`：41声明于 `include/pto/common/pto_instr.hpp`：
42 42 
43```cpp43```cpp
44-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>44+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
45+ typename TileDataSrc1, typename... WaitEvents>
45PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events);46PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events);
46```47```
47 48 
49+`PrecisionType`可指定以下值：
50+ 
51+* `DivAlgorithm::DEFAULT`：普通算法，速度快但精度较低。
52+* `DivAlgorithm::HIGH_PRECISION`：高精度算法，速度较慢。
53+ 
48## 约束54## 约束
49 55 
50- **实现检查 (A2A3)**:56- **实现检查 (A2A3)**:
@@ -63,6 +69,8 @@ PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &sr
63 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域;.69 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域;.
64- **除零**:70- **除零**:
65 - 行为由目标定义。71 - 行为由目标定义。
72+- **高精度算法**
73+ - 仅在A5上有效，`PrecisionType`选项A3上将被忽略。
66 74 
67## 示例75## 示例
68 76 
@@ -77,6 +85,7 @@ void example_auto() {
77 using TileT = Tile<TileType::Vec, float, 16, 16>;85 using TileT = Tile<TileType::Vec, float, 16, 16>;
78 TileT src0, src1, dst;86 TileT src0, src1, dst;
79 TDIV(dst, src0, src1);87 TDIV(dst, src0, src1);
88+ TDIV<DivAlgorithm::HIGH_PRECISION>(dst, src0, src1); // A5 Only
80}89}
81```90```
82 91 
@@ -94,6 +103,7 @@ void example_manual() {
94 TASSIGN(src1, 0x2000);103 TASSIGN(src1, 0x2000);
95 TASSIGN(dst, 0x3000);104 TASSIGN(dst, 0x3000);
96 TDIV(dst, src0, src1);105 TDIV(dst, src0, src1);
106+ TDIV<DivAlgorithm::HIGH_PRECISION>(dst, src0, src1); // A5 Only
97}107}
98```108```
99 109 
Mdocs/isa/TRECIP.md+10-1
@@ -41,10 +41,16 @@ pto.trecip ins(%src : !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
41Declared in `include/pto/common/pto_instr.hpp`:41Declared in `include/pto/common/pto_instr.hpp`:
42 42 
43```cpp43```cpp
44-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>44+template <auto PrecisionType = RecipAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
45+ typename... WaitEvents>
45PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events);46PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events);
46```47```
47 48 
49+`PrecisionType` has the following values available:
50+ 
51+* `RecipAlgorithm::DEFAULT`: Normal algorithm, faster but with lower precision.
52+* `RecipAlgorithm::HIGH_PRECISION`: High precision algorithm, but slower.
53+ 
48## Constraints54## Constraints
49 55 
50- **Implementation checks (NPU)**:56- **Implementation checks (NPU)**:
@@ -58,6 +64,8 @@ PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &...
58 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain.64 - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain.
59- **Domain / NaN**:65- **Domain / NaN**:
60 - Division-by-zero behavior is target-defined; the CPU simulator asserts in debug builds.66 - Division-by-zero behavior is target-defined; the CPU simulator asserts in debug builds.
67+- **High Precision Algorithm**
68+ - Only available on A5, `PrecisionType` option is ignored on A3.
61 69 
62## Examples70## Examples
63 71 
@@ -70,6 +78,7 @@ void example() {
70 using TileT = Tile<TileType::Vec, float, 16, 16>;78 using TileT = Tile<TileType::Vec, float, 16, 16>;
71 TileT x, out;79 TileT x, out;
72 TRECIP(out, x);80 TRECIP(out, x);
81+ TRECIP<RecipAlgorithm::HIGH_PRECISION>(out, x);
73}82}
74```83```
75 84 
Mdocs/isa/TRECIP_zh.md+10-1
@@ -41,10 +41,16 @@ pto.trecip ins(%src : !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
41声明于 `include/pto/common/pto_instr.hpp`：41声明于 `include/pto/common/pto_instr.hpp`：
42 42 
43```cpp43```cpp
44-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>44+template <auto PrecisionType = RecipAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
45+ typename... WaitEvents>
45PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events);46PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events);
46```47```
47 48 
49+`PrecisionType`可指定以下值：
50+ 
51+* `RecipAlgorithm::DEFAULT`：普通算法，速度快但精度较低。
52+* `RecipAlgorithm::HIGH_PRECISION`：高精度算法，速度较慢。
53+ 
48## 约束54## 约束
49 55 
50- **实现检查 (NPU)**:56- **实现检查 (NPU)**:
@@ -58,6 +64,8 @@ PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &...
58 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域。64 - 该操作使用 `dst.GetValidRow()` / `dst.GetValidCol()` 作为迭代域。
59- **域 / NaN**:65- **域 / NaN**:
60 - 除零行为由目标定义；CPU 模拟器在调试构建中会断言。66 - 除零行为由目标定义；CPU 模拟器在调试构建中会断言。
67+- **高精度算法**
68+ - 仅在A5上有效，`PrecisionType`选项A3上将被忽略。
61 69 
62## 示例70## 示例
63 71 
@@ -70,6 +78,7 @@ void example() {
70 using TileT = Tile<TileType::Vec, float, 16, 16>;78 using TileT = Tile<TileType::Vec, float, 16, 16>;
71 TileT x, out;79 TileT x, out;
72 TRECIP(out, x);80 TRECIP(out, x);
81+ TRECIP<RecipAlgorithm::HIGH_PRECISION>(out, x);
73}82}
74```83```
75 84 
Minclude/pto/common/pto_instr.hpp+12-8
@@ -290,11 +290,12 @@ PTO_INST RecordEvent TSTORE_FP(GlobalData &dst, TileData &src, FpTileData &fp, W
290 return {};290 return {};
291}291}
292 292 
293-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>293+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
294+ typename TileDataSrc1, typename... WaitEvents>
294PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events)295PTO_INST RecordEvent TDIV(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &... events)
295{296{
296 TSYNC(events...);297 TSYNC(events...);
297- MAP_INSTR_IMPL(TDIV, dst, src0, src1);298+ TDIV_IMPL<PrecisionType>(dst, src0, src1);
298 return {};299 return {};
299}300}
300 301 
@@ -348,14 +349,15 @@ PTO_INST RecordEvent TLOG(TileDataDst &dst, TileDataSrc &src, WaitEvents &... ev
348 return {};349 return {};
349}350}
350 351 
351-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>352+template <auto PrecisionType = RecipAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
353+ typename... WaitEvents>
352PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events)354PTO_INST RecordEvent TRECIP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events)
353{355{
354 TSYNC(events...);356 TSYNC(events...);
355 /*357 /*
356 * A3's TRECIP instruction does not support setting the source Tile and destination Tile to the same memory.358 * A3's TRECIP instruction does not support setting the source Tile and destination Tile to the same memory.
357 */359 */
358- MAP_INSTR_IMPL(TDIVS, dst, 1, src);360+ TDIVS_IMPL<static_cast<DivAlgorithm>(PrecisionType)>(dst, 1, src);
359 return {};361 return {};
360}362}
361 363 
@@ -1402,12 +1404,13 @@ PTO_INST RecordEvent TSUBS(TileDataDst &dst, TileDataSrc &src0, typename TileDat
1402 return {};1404 return {};
1403}1405}
1404 1406 
1405-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>1407+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
1408+ typename... WaitEvents>
1406PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar,1409PTO_INST RecordEvent TDIVS(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar,
1407 WaitEvents &... events)1410 WaitEvents &... events)
1408{1411{
1409 TSYNC(events...);1412 TSYNC(events...);
1410- MAP_INSTR_IMPL(TDIVS, dst, src0, scalar);1413+ TDIVS_IMPL<PrecisionType>(dst, src0, scalar);
1411 return {};1414 return {};
1412}1415}
1413 1416 
@@ -1420,12 +1423,13 @@ PTO_INST RecordEvent TMULS(TileDataDst &dst, TileDataSrc &src0, typename TileDat
1420 return {};1423 return {};
1421}1424}
1422 1425 
1423-template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>1426+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
1427+ typename... WaitEvents>
1424PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0,1428PTO_INST RecordEvent TDIVS(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src0,
1425 WaitEvents &... events)1429 WaitEvents &... events)
1426{1430{
1427 TSYNC(events...);1431 TSYNC(events...);
1428- MAP_INSTR_IMPL(TDIVS, dst, scalar, src0);1432+ TDIVS_IMPL<PrecisionType>(dst, scalar, src0);
1429 return {};1433 return {};
1430}1434}
1431 1435 
Minclude/pto/common/type.hpp+12-0
@@ -294,6 +294,18 @@ enum class TileLayoutCustom : uint8_t
294 NONE,294 NONE,
295};295};
296 296 
297+enum class DivAlgorithm : uint8_t
298+{
299+ DEFAULT,
300+ HIGH_PRECISION
301+};
302+ 
303+enum class RecipAlgorithm : uint8_t
304+{
305+ DEFAULT,
306+ HIGH_PRECISION
307+};
308+ 
297namespace GlobalTensorDim {309namespace GlobalTensorDim {
298constexpr int DIM_0 = 0;310constexpr int DIM_0 = 0;
299constexpr int DIM_1 = 1;311constexpr int DIM_1 = 1;
Minclude/pto/costmodel/a2a3/TDivS.hpp+2-2
@@ -14,13 +14,13 @@ See LICENSE in the root of the software repository for the full text of the Lice
14 14 
15namespace pto {15namespace pto {
16 16 
17-template <typename TileDataDst, typename TileDataSrc>17+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
18PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src, typename TileDataSrc::DType scalar)18PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src, typename TileDataSrc::DType scalar)
19{19{
20 pto::CostModel::GetInstance().BinSOpPredictCycle<DivSOp, TileDataDst, TileDataSrc>("TDIVS", dst, src);20 pto::CostModel::GetInstance().BinSOpPredictCycle<DivSOp, TileDataDst, TileDataSrc>("TDIVS", dst, src);
21}21}
22 22 
23-template <typename TileDataDst, typename TileDataSrc>23+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
24PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataSrc::DType scalar, TileDataSrc &src)24PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataSrc::DType scalar, TileDataSrc &src)
25{25{
26 pto::CostModel::GetInstance().BinSOpPredictCycle<DivSOp, TileDataDst, TileDataSrc>("TDIVS", dst, src);26 pto::CostModel::GetInstance().BinSOpPredictCycle<DivSOp, TileDataDst, TileDataSrc>("TDIVS", dst, src);
Minclude/pto/cpu/ElementTileOp.h+6-1
@@ -114,10 +114,15 @@ BINARY_OP_DEF(MIN)
114UNARY_OP_DEF(LOG)114UNARY_OP_DEF(LOG)
115UNARY_OP_DEF(NEG)115UNARY_OP_DEF(NEG)
116UNARY_OP_DEF(NOT)116UNARY_OP_DEF(NOT)
117-UNARY_OP_DEF(RECIP)
118UNARY_OP_DEF(RELU)117UNARY_OP_DEF(RELU)
119UNARY_OP_DEF(EXP)118UNARY_OP_DEF(EXP)
120 119 
120+template <auto PrecisionType = RecipAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
121+PTO_INTERNAL void TRECIP_IMPL(TileDataDst &dst, TileDataSrc &src)
122+{
123+ UnaryElementTileOp_Impl<ElementOp::OP_RECIP>(dst, src);
124+}
125+ 
121template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename TileDataTmp>126template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename TileDataTmp>
122PTO_INTERNAL void TREM_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, TileDataTmp &tmp)127PTO_INTERNAL void TREM_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, TileDataTmp &tmp)
123{128{
Minclude/pto/cpu/TBinSOps.hpp+2-2
@@ -45,7 +45,7 @@ PTO_INTERNAL void TMULS_IMPL(TileData &dst, TileData &src, typename TileData::DT
45 unsigned col = dst.GetValidCol();45 unsigned col = dst.GetValidCol();
46 TBinSOp<TileData>(dst.data(), src.data(), scalar, row, col, lambda);46 TBinSOp<TileData>(dst.data(), src.data(), scalar, row, col, lambda);
47}47}
48-template <typename TileData>48+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileData>
49PTO_INTERNAL void TDIVS_IMPL(TileData &dst, TileData &src, typename TileData::DType scalar)49PTO_INTERNAL void TDIVS_IMPL(TileData &dst, TileData &src, typename TileData::DType scalar)
50{50{
51 auto lambda = [](typename TileData::DType x, typename TileData::DType y) { return x / y; };51 auto lambda = [](typename TileData::DType x, typename TileData::DType y) { return x / y; };
@@ -53,7 +53,7 @@ PTO_INTERNAL void TDIVS_IMPL(TileData &dst, TileData &src, typename TileData::DT
53 unsigned col = dst.GetValidCol();53 unsigned col = dst.GetValidCol();
54 TBinSOp<TileData>(dst.data(), src.data(), scalar, row, col, lambda);54 TBinSOp<TileData>(dst.data(), src.data(), scalar, row, col, lambda);
55}55}
56-template <typename TileData>56+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileData>
57PTO_INTERNAL void TDIVS_IMPL(TileData &dst, typename TileData::DType scalar, TileData &src)57PTO_INTERNAL void TDIVS_IMPL(TileData &dst, typename TileData::DType scalar, TileData &src)
58{58{
59 auto lambda = [](typename TileData::DType x, typename TileData::DType y) { return y / x; };59 auto lambda = [](typename TileData::DType x, typename TileData::DType y) { return y / x; };
Minclude/pto/cpu/TDiv.hpp+1-1
@@ -59,7 +59,7 @@ void TDiv_Impl(typename tile_shape::TileDType dst, typename tile_shape::TileDTyp
59 }59 }
60}60}
61 61 
62-template <typename tile_shape>62+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename tile_shape>
63PTO_INTERNAL void TDIV_IMPL(tile_shape &dst, tile_shape &src0, tile_shape &src1)63PTO_INTERNAL void TDIV_IMPL(tile_shape &dst, tile_shape &src0, tile_shape &src1)
64{64{
65 unsigned row = dst.GetValidRow();65 unsigned row = dst.GetValidRow();
Minclude/pto/npu/a2a3/TDiv.hpp+2-1
@@ -65,7 +65,8 @@ PTO_INTERNAL void TDivCheck(const TileDataDst &dst, const TileDataSrc0 &src0, co
65 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");65 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");
66}66}
67 67 
68-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>68+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
69+ typename TileDataSrc1>
69PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)70PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)
70{71{
71 using T = typename TileDataDst::DType;72 using T = typename TileDataDst::DType;
Minclude/pto/npu/a2a3/TDivS.hpp+2-2
@@ -201,7 +201,7 @@ __tf__ PTO_INTERNAL void TDivS(typename TileDataDst::TileDType __out__ dstData,
201 }201 }
202}202}
203 203 
204-template <typename TileDataDst, typename TileDataSrc>204+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
205PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src, typename TileDataSrc::DType scalar)205PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src, typename TileDataSrc::DType scalar)
206{206{
207 using T = typename TileDataSrc::DType;207 using T = typename TileDataSrc::DType;
@@ -254,7 +254,7 @@ __tf__ PTO_INTERNAL void TSDiv(typename TileDataDst::TileDType __out__ dstData,
254 dst, src, scalar, validRow, validCol);254 dst, src, scalar, validRow, validCol);
255 }255 }
256}256}
257-template <typename TileDataDst, typename TileDataSrc>257+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
258PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src)258PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataDst::DType scalar, TileDataSrc &src)
259{259{
260 using T = typename TileDataSrc::DType;260 using T = typename TileDataSrc::DType;
Minclude/pto/npu/a5/TDiv.hpp+10-8
@@ -21,9 +21,10 @@ See LICENSE in the root of the software repository for the full text of the Lice
21#ifndef STRAIGHT_INTRINSICS_IMPL21#ifndef STRAIGHT_INTRINSICS_IMPL
22#include "custom/Div754.hpp"22#include "custom/Div754.hpp"
23#endif23#endif
24+ 
24namespace pto {25namespace pto {
25 26 
26-template <typename T>27+template <DivAlgorithm PrecisionType, typename T>
27struct DivOp {28struct DivOp {
28#ifdef STRAIGHT_INTRINSICS_IMPL29#ifdef STRAIGHT_INTRINSICS_IMPL
29 PTO_INTERNAL static void BinInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, RegTensor<T> &reg_src1,30 PTO_INTERNAL static void BinInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, RegTensor<T> &reg_src1,
@@ -35,9 +36,9 @@ struct DivOp {
35 PTO_INTERNAL static void BinInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, RegTensor<T> &reg_src1,36 PTO_INTERNAL static void BinInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, RegTensor<T> &reg_src1,
36 MaskReg &preg)37 MaskReg &preg)
37 {38 {
38- if constexpr (std::is_same_v<T, float>) {39+ if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, float>) {
39 DivIEEE754FloatImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_src1, preg);40 DivIEEE754FloatImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_src1, preg);
40- } else if constexpr (std::is_same_v<T, half>) {41+ } else if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, half>) {
41 DivIEEE754HalfImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_src1, preg);42 DivIEEE754HalfImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_src1, preg);
42 } else {43 } else {
43 vdiv(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING);44 vdiv(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING);
@@ -46,8 +47,8 @@ struct DivOp {
46#endif47#endif
47};48};
48 49 
49-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, unsigned ElementsPerRepeat,50+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
50- unsigned BlockSizeElem>51+ typename TileDataSrc1, unsigned ElementsPerRepeat, unsigned BlockSizeElem>
51__tf__ PTO_INTERNAL OP_NAME(TDIV)52__tf__ PTO_INTERNAL OP_NAME(TDIV)
52 OP_TYPE(element_wise) void TDiv(typename TileDataDst::TileDType __out__ dst,53 OP_TYPE(element_wise) void TDiv(typename TileDataDst::TileDType __out__ dst,
53 typename TileDataSrc0::TileDType __in__ src0,54 typename TileDataSrc0::TileDType __in__ src0,
@@ -59,7 +60,7 @@ __tf__ PTO_INTERNAL OP_NAME(TDIV)
59 __ubuf__ T *src0Ptr = (__ubuf__ T *)__cce_get_tile_ptr(src0);60 __ubuf__ T *src0Ptr = (__ubuf__ T *)__cce_get_tile_ptr(src0);
60 __ubuf__ T *src1Ptr = (__ubuf__ T *)__cce_get_tile_ptr(src1);61 __ubuf__ T *src1Ptr = (__ubuf__ T *)__cce_get_tile_ptr(src1);
61 62 
62- BinaryInstr<DivOp<T>, TileDataDst, TileDataSrc0, TileDataSrc1, ElementsPerRepeat, BlockSizeElem>(63+ BinaryInstr<DivOp<PrecisionType, T>, TileDataDst, TileDataSrc0, TileDataSrc1, ElementsPerRepeat, BlockSizeElem>(
63 dstPtr, src0Ptr, src1Ptr, validRows, validCols, version);64 dstPtr, src0Ptr, src1Ptr, validRows, validCols, version);
64 return;65 return;
65}66}
@@ -84,7 +85,8 @@ PTO_INTERNAL void TDivCheck(const TileDataDst &dst, const TileDataSrc0 &src0, co
84 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");85 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");
85}86}
86 87 
87-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>88+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
89+ typename TileDataSrc1>
88PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)90PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)
89{91{
90 using T = typename TileDataDst::DType;92 using T = typename TileDataDst::DType;
@@ -92,7 +94,7 @@ PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &
92 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);94 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);
93 constexpr unsigned elementsPerRepeat = REPEAT_BYTE / sizeof(T);95 constexpr unsigned elementsPerRepeat = REPEAT_BYTE / sizeof(T);
94 96 
95- TDiv<TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem>(97+ TDiv<PrecisionType, TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem>(
96 dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());98 dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());
97}99}
98} // namespace pto100} // namespace pto
Minclude/pto/npu/a5/TDivS.hpp+25-20
@@ -22,24 +22,23 @@ See LICENSE in the root of the software repository for the full text of the Lice
22 22 
23namespace pto {23namespace pto {
24 24 
25-template <typename T>25+template <DivAlgorithm PrecisionType, typename T>
26struct DivSOp {26struct DivSOp {
27#ifndef STRAIGHT_INTRINSICS_IMPL27#ifndef STRAIGHT_INTRINSICS_IMPL
28 PTO_INTERNAL static void BinSInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, T reg_src1, MaskReg &preg)28 PTO_INTERNAL static void BinSInstr(RegTensor<T> &reg_dst, RegTensor<T> &reg_src0, T reg_src1, MaskReg &preg)
29 {29 {
30- if constexpr (std::is_same_v<T, float>) {30+ if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, float>) {
31 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);31 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);
32- DivIEEE754FloatImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_dst, preg);32+ DivIEEE754FloatImpl<T, RegTensor<T>>(reg_dst, reg_src0, reg_dst, preg);
33- } else if constexpr (std::is_same_v<T, half>) {33+ } else if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, half>) {
34 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);34 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);
35- DivIEEE754HalfImpl<T, RegTensor<T> >(reg_dst, reg_src0, reg_dst, preg);35+ DivIEEE754HalfImpl<T, RegTensor<T>>(reg_dst, reg_src0, reg_dst, preg);
36 } else {36 } else {
37 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);37 vdup(reg_dst, reg_src1, preg, MODE_ZEROING);
38 vdiv(reg_dst, reg_src0, reg_dst, preg, MODE_ZEROING);38 vdiv(reg_dst, reg_src0, reg_dst, preg, MODE_ZEROING);
39 }39 }
40 }40 }
41#else41#else
42- 
43 PTO_INTERNAL static void BinSInstr(RegTensor<T> &vregdst, RegTensor<T> &vregsrc, T src1, MaskReg &preg)42 PTO_INTERNAL static void BinSInstr(RegTensor<T> &vregdst, RegTensor<T> &vregsrc, T src1, MaskReg &preg)
44 {43 {
45 float divider = static_cast<float>(src1);44 float divider = static_cast<float>(src1);
@@ -69,11 +68,17 @@ struct DivSOp {
69#endif68#endif
70};69};
71 70 
72-template <typename T>71+template <enum DivAlgorithm PrecisionType, typename T>
73struct DivSOpS {72struct DivSOpS {
74 PTO_INTERNAL static void BinSInstr(RegTensor<T> &vregdst, RegTensor<T> &vregsrc, T src0, MaskReg &preg)73 PTO_INTERNAL static void BinSInstr(RegTensor<T> &vregdst, RegTensor<T> &vregsrc, T src0, MaskReg &preg)
75 {74 {
76- if constexpr (std::is_same<T, float>::value || std::is_same<T, half>::value) {75+ if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, float>) {
76+ vdup(vregdst, src0, preg, MODE_ZEROING);
77+ DivIEEE754FloatImpl<T, RegTensor<T>>(vregdst, vregdst, vregsrc, preg);
78+ } else if constexpr (PrecisionType == DivAlgorithm::HIGH_PRECISION && std::is_same_v<T, half>) {
79+ vdup(vregdst, src0, preg, MODE_ZEROING);
80+ DivIEEE754HalfImpl<T, RegTensor<T>>(vregdst, vregdst, vregsrc, preg);
81+ } else if constexpr (std::is_same<T, float>::value || std::is_same<T, half>::value) {
77 vdup(vregdst, src0, preg, MODE_ZEROING);82 vdup(vregdst, src0, preg, MODE_ZEROING);
78 vdiv(vregdst, vregdst, vregsrc, preg);83 vdiv(vregdst, vregdst, vregsrc, preg);
79 } else if constexpr (std::is_same<T, int32_t>::value) {84 } else if constexpr (std::is_same<T, int32_t>::value) {
@@ -126,8 +131,8 @@ PTO_INTERNAL void TSDiv_naive(__ubuf__ T *dst, __ubuf__ T *src0, T src1, unsigne
126 }131 }
127 }132 }
128}133}
129-template <typename TileDataDst, typename TileDataSrc, unsigned elementsPerRepeat, unsigned blockSizeElem,134+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
130- unsigned dstRowStride, unsigned srcRowStride>135+ unsigned elementsPerRepeat, unsigned blockSizeElem, unsigned dstRowStride, unsigned srcRowStride>
131__tf__ PTO_INTERNAL OP_NAME(TDIVS)136__tf__ PTO_INTERNAL OP_NAME(TDIVS)
132 OP_TYPE(element_wise) void TDivS(typename TileDataDst::TileDType __out__ dst,137 OP_TYPE(element_wise) void TDivS(typename TileDataDst::TileDType __out__ dst,
133 typename TileDataSrc::TileDType __in__ src0,138 typename TileDataSrc::TileDType __in__ src0,
@@ -140,13 +145,13 @@ __tf__ PTO_INTERNAL OP_NAME(TDIVS)
140 if constexpr (std::is_integral_v<T>) {145 if constexpr (std::is_integral_v<T>) {
141 TDivs_naive<T, TileDataDst::Cols, TileDataSrc::Cols>(dstPtr, src0Ptr, src1, validRow, validCol);146 TDivs_naive<T, TileDataDst::Cols, TileDataSrc::Cols>(dstPtr, src0Ptr, src1, validRow, validCol);
142 } else {147 } else {
143- BinaryInstr<DivSOp<T>, TileDataDst, TileDataSrc, T, elementsPerRepeat, blockSizeElem, dstRowStride,148+ BinaryInstr<DivSOp<PrecisionType, T>, TileDataDst, TileDataSrc, T, elementsPerRepeat, blockSizeElem,
144- srcRowStride>(dstPtr, src0Ptr, src1, validRow, validCol, version);149+ dstRowStride, srcRowStride>(dstPtr, src0Ptr, src1, validRow, validCol, version);
145 }150 }
146}151}
147 152 
148-template <typename TileDataDst, typename TileDataSrc, unsigned elementsPerRepeat, unsigned blockSizeElem,153+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc,
149- unsigned dstRowStride, unsigned srcRowStride>154+ unsigned elementsPerRepeat, unsigned blockSizeElem, unsigned dstRowStride, unsigned srcRowStride>
150__tf__ PTO_INTERNAL OP_NAME(TDIVS)155__tf__ PTO_INTERNAL OP_NAME(TDIVS)
151 OP_TYPE(element_wise) void TDivS(typename TileDataDst::TileDType __out__ dst,156 OP_TYPE(element_wise) void TDivS(typename TileDataDst::TileDType __out__ dst,
152 typename TileDataSrc::DType __in__ src1,157 typename TileDataSrc::DType __in__ src1,
@@ -159,12 +164,12 @@ __tf__ PTO_INTERNAL OP_NAME(TDIVS)
159 if constexpr (std::is_integral_v<T>) {164 if constexpr (std::is_integral_v<T>) {
160 TSDiv_naive<T, TileDataDst::Cols, TileDataSrc::Cols>(dstPtr, src0Ptr, src1, validRow, validCol);165 TSDiv_naive<T, TileDataDst::Cols, TileDataSrc::Cols>(dstPtr, src0Ptr, src1, validRow, validCol);
161 } else {166 } else {
162- BinaryInstr<DivSOpS<T>, TileDataDst, TileDataSrc, T, elementsPerRepeat, blockSizeElem, dstRowStride,167+ BinaryInstr<DivSOpS<PrecisionType, T>, TileDataDst, TileDataSrc, T, elementsPerRepeat, blockSizeElem,
163- srcRowStride>(dstPtr, src0Ptr, src1, validRow, validCol, version);168+ dstRowStride, srcRowStride>(dstPtr, src0Ptr, src1, validRow, validCol, version);
164 }169 }
165}170}
166 171 
167-template <typename TileDataDst, typename TileDataSrc>172+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
168PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar)173PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src0, typename TileDataSrc::DType scalar)
169{174{
170 static_assert(std::is_same<typename TileDataDst::DType, uint32_t>::value ||175 static_assert(std::is_same<typename TileDataDst::DType, uint32_t>::value ||
@@ -200,11 +205,11 @@ PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, TileDataSrc &src0, typename TileD
200 constexpr unsigned srcRowStride = TileDataSrc::RowStride;205 constexpr unsigned srcRowStride = TileDataSrc::RowStride;
201 unsigned validRow = dst.GetValidRow();206 unsigned validRow = dst.GetValidRow();
202 unsigned validCol = dst.GetValidCol();207 unsigned validCol = dst.GetValidCol();
203- TDivS<TileDataDst, TileDataSrc, elementsPerRepeat, blockSizeElem, dstRowStride, srcRowStride>(208+ TDivS<PrecisionType, TileDataDst, TileDataSrc, elementsPerRepeat, blockSizeElem, dstRowStride, srcRowStride>(
204 dst.data(), src0.data(), scalar, validRow, validCol);209 dst.data(), src0.data(), scalar, validRow, validCol);
205}210}
206 211 
207-template <typename TileDataDst, typename TileDataSrc>212+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc>
208PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataSrc::DType scalar, TileDataSrc &src0)213PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataSrc::DType scalar, TileDataSrc &src0)
209{214{
210 static_assert(TileDataSrc::Loc == TileType::Vec, "TileType of src and dst tiles must be TileType::Vec.");215 static_assert(TileDataSrc::Loc == TileType::Vec, "TileType of src and dst tiles must be TileType::Vec.");
@@ -227,7 +232,7 @@ PTO_INTERNAL void TDIVS_IMPL(TileDataDst &dst, typename TileDataSrc::DType scala
227 constexpr unsigned srcRowStride = TileDataSrc::RowStride;232 constexpr unsigned srcRowStride = TileDataSrc::RowStride;
228 unsigned validRow = dst.GetValidRow();233 unsigned validRow = dst.GetValidRow();
229 unsigned validCol = dst.GetValidCol();234 unsigned validCol = dst.GetValidCol();
230- TDivS<TileDataDst, TileDataSrc, elementsPerRepeat, blockSizeElem, dstRowStride, srcRowStride>(235+ TDivS<PrecisionType, TileDataDst, TileDataSrc, elementsPerRepeat, blockSizeElem, dstRowStride, srcRowStride>(
231 dst.data(), scalar, src0.data(), validRow, validCol);236 dst.data(), scalar, src0.data(), validRow, validCol);
232}237}
233} // namespace pto238} // namespace pto
Minclude/pto/npu/kirin9030/TDiv.hpp+2-1
@@ -67,7 +67,8 @@ PTO_INTERNAL void TDivCheck(const TileDataDst &dst, const TileDataSrc0 &src0, co
67 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");67 "Fix: TDIV input tile src1 valid shape mismatch with output tile dst shape.");
68}68}
69 69 
70-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>70+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc0,
71+ typename TileDataSrc1>
71PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)72PTO_INTERNAL void TDIV_IMPL(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1)
72{73{
73 using T = typename TileDataDst::DType;74 using T = typename TileDataDst::DType;
Minclude/pto/npu/kirin9030/TDivS.hpp+2-2
@@ -63,7 +63,7 @@ __tf__ PTO_INTERNAL OP_NAME(TDIVS)
63 dstPtr, src0Ptr, src1, validRow, validCol, version);63 dstPtr, src0Ptr, src1, validRow, validCol, version);
64}64}
65 65 
66-template <typename DstTile, typename SrcTile>66+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename DstTile, typename SrcTile>
67PTO_INTERNAL void TDIVS_IMPL(DstTile &dst, SrcTile &src0, typename SrcTile::DType scalar)67PTO_INTERNAL void TDIVS_IMPL(DstTile &dst, SrcTile &src0, typename SrcTile::DType scalar)
68{68{
69 static_assert(std::is_same<typename DstTile::DType, uint32_t>::value ||69 static_assert(std::is_same<typename DstTile::DType, uint32_t>::value ||
@@ -103,7 +103,7 @@ PTO_INTERNAL void TDIVS_IMPL(DstTile &dst, SrcTile &src0, typename SrcTile::DTyp
103 scalar, validRow, validCol);103 scalar, validRow, validCol);
104}104}
105 105 
106-template <typename DstTile, typename SrcTile>106+template <auto PrecisionType = DivAlgorithm::DEFAULT, typename DstTile, typename SrcTile>
107PTO_INTERNAL void TDIVS_IMPL(DstTile &dst, typename SrcTile::DType scalar, SrcTile &src0)107PTO_INTERNAL void TDIVS_IMPL(DstTile &dst, typename SrcTile::DType scalar, SrcTile &src0)
108{108{
109 static_assert(SrcTile::Loc == TileType::Vec, "TileType of src and dst tiles must be TileType::Vec.");109 static_assert(SrcTile::Loc == TileType::Vec, "TileType of src and dst tiles must be TileType::Vec.");
Mtests/npu/a5/src/st/testcase/tdiv/gen_data.py+27-12
@@ -12,7 +12,6 @@
12 12 
13import os13import os
14import numpy as np14import numpy as np
15-np.random.seed(19)
16 15 
17 16 
18def gen_golden_data(case_name, param):17def gen_golden_data(case_name, param):
@@ -24,8 +23,18 @@ def gen_golden_data(case_name, param):
24 h_valid, w_valid = param.valid_row, param.valid_col23 h_valid, w_valid = param.valid_row, param.valid_col
25 24 
26 # Generate random input arrays25 # Generate random input arrays
27- input1 = np.random.randint(1, 10, size=[src0_tile_row, src0_tile_col]).astype(dtype)26+ if dtype in (np.int8, np.uint8, np.int16, np.uint16, np.int32, np.uint32):
28- input2 = np.random.randint(1, 10, size=[src1_tile_row, src1_tile_col]).astype(dtype)27+ dtype_info = np.iinfo(dtype)
28+ input1 = np.random.randint(dtype_info.min, dtype_info.max,
29+ size=[src0_tile_row, src0_tile_col]).astype(dtype)
30+ input2 = np.random.randint(dtype_info.min, dtype_info.max,
31+ size=[src1_tile_row, src1_tile_col]).astype(dtype)
32+ else:
33+ dtype_info = np.finfo(dtype)
34+ input1 = np.random.uniform(low=dtype_info.min, high=dtype_info.max,
35+ size=[src0_tile_row, src0_tile_col]).astype(dtype)
36+ input2 = np.random.uniform(low=dtype_info.min, high=dtype_info.max,
37+ size=[src1_tile_row, src1_tile_col]).astype(dtype)
29 38 
30 # Perform the operation39 # Perform the operation
31 golden = np.zeros([dst_tile_row, dst_tile_col]).astype(dtype)40 golden = np.zeros([dst_tile_row, dst_tile_col]).astype(dtype)
@@ -38,16 +47,18 @@ def gen_golden_data(case_name, param):
38 47 
39 48 
40class TDivParams:49class TDivParams:
41- def __init__(self, dtype, dstH, dstW, src0H, src0W, src1H, src1W, vRow, vCol):50+ def __init__(self, dtype, dst_tile_row, dst_tile_col, src0_tile_row, src0_tile_col,
51+ src1_tile_row, src1_tile_col, valid_row, valid_col, high_precision=False):
42 self.dtype = dtype52 self.dtype = dtype
43- self.dst_tile_row = dstH53+ self.dst_tile_row = dst_tile_row
44- self.dst_tile_col = dstW54+ self.dst_tile_col = dst_tile_col
45- self.src0_tile_row = src0H55+ self.src0_tile_row = src0_tile_row
46- self.src0_tile_col = src0W56+ self.src0_tile_col = src0_tile_col
47- self.src1_tile_row = src1H57+ self.src1_tile_row = src1_tile_row
48- self.src1_tile_col = src1W58+ self.src1_tile_col = src1_tile_col
49- self.valid_row = vRow59+ self.valid_row = valid_row
50- self.valid_col = vCol60+ self.valid_col = valid_col
61+ self.high_precision = high_precision
51 62 
52 63 
53def generate_case_name(param):64def generate_case_name(param):
@@ -58,6 +69,8 @@ def generate_case_name(param):
58 np.int32: 'int32',69 np.int32: 'int32',
59 np.int16: 'int16'70 np.int16: 'int16'
60 }[param.dtype]71 }[param.dtype]
72+ if param.high_precision:
73+ dtype_str += '_hp'
61 return f"TDIVTest.case_{dtype_str}_{param.dst_tile_row}x{param.dst_tile_col}_\74 return f"TDIVTest.case_{dtype_str}_{param.dst_tile_row}x{param.dst_tile_col}_\
62{param.src0_tile_row}x{param.src0_tile_col}_{param.src1_tile_row}x{param.src1_tile_col}_\75{param.src0_tile_row}x{param.src0_tile_col}_{param.src1_tile_row}x{param.src1_tile_col}_\
63{param.valid_row}x{param.valid_col}"76{param.valid_row}x{param.valid_col}"
@@ -84,6 +97,8 @@ if __name__ == "__main__":
84 TDivParams(np.float32, 16, 32, 16, 64, 16, 32, 16, 31),97 TDivParams(np.float32, 16, 32, 16, 64, 16, 32, 16, 31),
85 TDivParams(np.int16, 32, 128, 32, 128, 32, 256, 32, 127),98 TDivParams(np.int16, 32, 128, 32, 128, 32, 256, 32, 127),
86 TDivParams(np.int32, 16, 32, 16, 64, 16, 32, 16, 31),99 TDivParams(np.int32, 16, 32, 16, 64, 16, 32, 16, 31),
100+ TDivParams(np.float32, 2, 16, 2, 16, 2, 16, 2, 16, True),
101+ TDivParams(np.float16, 2, 32, 2, 32, 2, 32, 2, 32, True),
87 ]102 ]
88 103 
89 for param in case_params_list:104 for param in case_params_list:
Mtests/npu/a5/src/st/testcase/tdiv/main.cpp+15-8
@@ -33,17 +33,15 @@ std::string GetGoldenDir()
33}33}
34 34 
35template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,35template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
36- int vCols, bool sameTile>36+ int vCols, bool highPrecision>
37void LaunchTDiv(T *out, T *src0, T *src1, void *stream);37void LaunchTDiv(T *out, T *src0, T *src1, void *stream);
38 38 
39template <int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows, int vCols,39template <int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows, int vCols,
40- bool sameTile>40+ bool highPrecision>
41void LaunchTDivHalf(aclFloat16 *out, aclFloat16 *src0, aclFloat16 *src1, void *stream);41void LaunchTDivHalf(aclFloat16 *out, aclFloat16 *src0, aclFloat16 *src1, void *stream);
42 42 
43template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,43template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
44- int vCols, bool isHalf = false,44+ int vCols, bool isHalf = false, bool highPrecision = false>
45- bool sameTile = (dstTileH == src0TileH && dstTileH == src1TileH && dstTileW == src0TileW &&
46- dstTileW == src1TileW)>
47void test_tdiv()45void test_tdiv()
48{46{
49 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);47 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);
@@ -72,10 +70,10 @@ void test_tdiv()
72 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);70 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);
73 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);71 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);
74 if constexpr (isHalf) {72 if constexpr (isHalf) {
75- LaunchTDivHalf<dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols, sameTile>(73+ LaunchTDivHalf<dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols, highPrecision>(
76 dstDevice, src0Device, src1Device, stream);74 dstDevice, src0Device, src1Device, stream);
77 } else {75 } else {
78- LaunchTDiv<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols, sameTile>(76+ LaunchTDiv<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols, highPrecision>(
79 dstDevice, src0Device, src1Device, stream);77 dstDevice, src0Device, src1Device, stream);
80 }78 }
81 79 
@@ -100,7 +98,8 @@ void test_tdiv()
100 ReadFile(GetGoldenDir() + "/golden.bin", fileSizeDst, golden.data(), fileSizeDst);98 ReadFile(GetGoldenDir() + "/golden.bin", fileSizeDst, golden.data(), fileSizeDst);
101 ReadFile(GetGoldenDir() + "/output.bin", fileSizeDst, devFinal.data(), fileSizeDst);99 ReadFile(GetGoldenDir() + "/output.bin", fileSizeDst, devFinal.data(), fileSizeDst);
102 100 
103- bool ret = ResultCmp<T>(golden, devFinal, 0.001f);101+ auto resPrecision = highPrecision ? 0.0000001f : 0.001f;
102+ bool ret = ResultCmp<T>(golden, devFinal, resPrecision);
104 103 
105 EXPECT_TRUE(ret);104 EXPECT_TRUE(ret);
106}105}
@@ -153,3 +152,11 @@ TEST_F(TDIVTest, case_int32_16x32_16x64_16x32_16x31)
153{152{
154 test_tdiv<int32_t, 16, 32, 16, 64, 16, 32, 16, 31>();153 test_tdiv<int32_t, 16, 32, 16, 64, 16, 32, 16, 31>();
155}154}
155+TEST_F(TDIVTest, case_float_hp_2x16_2x16_2x16_2x16)
156+{
157+ test_tdiv<float, 2, 16, 2, 16, 2, 16, 2, 16, false, true>();
158+}
159+TEST_F(TDIVTest, case_half_hp_2x32_2x32_2x32_2x32)
160+{
161+ test_tdiv<aclFloat16, 2, 32, 2, 32, 2, 32, 2, 32, true, true>();
162+}
Mtests/npu/a5/src/st/testcase/tdiv/tdiv_kernel.cpp+20-59
@@ -14,42 +14,8 @@ See LICENSE in the root of the software repository for the full text of the Lice
14 14 
15using namespace pto;15using namespace pto;
16 16 
17-template <typename T, int kTRows_, int kTCols_, int kGRows_, int kGCols_>
18-__global__ AICORE void runTDIV(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1)
19-{
20- using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>;
21- using DynStridDim5 = pto::Stride<1, 1, 1, kGCols_, 1>;
22- using GlobalData = GlobalTensor<T, DynShapeDim5, DynStridDim5>;
23- using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>;
24- TileData src0Tile(kTRows_, kTCols_);
25- TileData src1Tile(kTRows_, kTCols_);
26- TileData dstTile(kTRows_, kTCols_);
27- TASSIGN(src0Tile, 0x0 + 0x400 * block_idx);
28- TASSIGN(src1Tile, 0x4000 + 0x400 * block_idx);
29- TASSIGN(dstTile, 0x8000 + 0x400 * block_idx);
30- 
31- int offset = (block_idx / 4) * (64 * 16) + (block_idx % 4) * 16;
32- GlobalData src0Global(src0 + offset);
33- GlobalData src1Global(src1 + offset);
34- GlobalData dstGlobal(out + offset);
35- 
36- TLOAD(src0Tile, src0Global);
37- TLOAD(src1Tile, src1Global);
38-#ifndef __PTO_AUTO__
39- set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
40- wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
41-#endif
42- TDIV(dstTile, src0Tile, src1Tile);
43-#ifndef __PTO_AUTO__
44- set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
45- wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
46-#endif
47- TSTORE(dstGlobal, dstTile);
48- out = dstGlobal.data();
49-}
50- 
51template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,17template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
52- int vCols>18+ int vCols, bool highPrecision = false>
53__global__ AICORE void runTDIV(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1)19__global__ AICORE void runTDIV(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1)
54{20{
55 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;21 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;
@@ -80,7 +46,8 @@ __global__ AICORE void runTDIV(__gm__ T __out__ *out, __gm__ T __in__ *src0, __g
80 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);46 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
81 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);47 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
82#endif48#endif
83- TDIV<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);49+ constexpr auto precisionType = highPrecision ? DivAlgorithm::HIGH_PRECISION : DivAlgorithm::DEFAULT;
50+ TDIV<precisionType>(dstTile, src0Tile, src1Tile);
84#ifndef __PTO_AUTO__51#ifndef __PTO_AUTO__
85 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);52 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
86 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);53 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
@@ -90,38 +57,29 @@ __global__ AICORE void runTDIV(__gm__ T __out__ *out, __gm__ T __in__ *src0, __g
90}57}
91 58 
92template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,59template <typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
93- int vCols, bool sameTile>60+ int vCols, bool highPrecision>
94void LaunchTDiv(T *out, T *src0, T *src1, void *stream)61void LaunchTDiv(T *out, T *src0, T *src1, void *stream)
95{62{
96- if constexpr (sameTile) {63+ runTDIV<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
97- runTDIV<T, dstTileH, dstTileW, vRows, vCols><<<1, nullptr, stream>>>(out, src0, src1);64+ <<<1, nullptr, stream>>>(out, src0, src1);
98- } else {
99- runTDIV<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
100- <<<1, nullptr, stream>>>(out, src0, src1);
101- }
102}65}
103 66 
104template <int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows, int vCols,67template <int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows, int vCols,
105- bool sameTile>68+ bool highPrecision>
106void LaunchTDivHalf(aclFloat16 *out, aclFloat16 *src0, aclFloat16 *src1, void *stream)69void LaunchTDivHalf(aclFloat16 *out, aclFloat16 *src0, aclFloat16 *src1, void *stream)
107{70{
108- if constexpr (sameTile) {71+ runTDIV<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
109- runTDIV<half, dstTileH, dstTileW, vRows, vCols>72+ <<<1, nullptr, stream>>>((half *)(out), (half *)(src0), (half *)(src1));
110- <<<1, nullptr, stream>>>((half *)(out), (half *)(src0), (half *)(src1));
111- } else {
112- runTDIV<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
113- <<<1, nullptr, stream>>>((half *)(out), (half *)(src0), (half *)(src1));
114- }
115}73}
116 74 
117-template void LaunchTDiv<float, 64, 64, 64, 64, 64, 64, 64, 64, true>(float *out, float *src0, float *src1,75+template void LaunchTDiv<float, 64, 64, 64, 64, 64, 64, 64, 64, false>(float *out, float *src0, float *src1,
118- void *stream);76+ void *stream);
119-template void LaunchTDiv<int32_t, 64, 64, 64, 64, 64, 64, 64, 64, true>(int32_t *out, int32_t *src0, int32_t *src1,77+template void LaunchTDiv<int32_t, 64, 64, 64, 64, 64, 64, 64, 64, false>(int32_t *out, int32_t *src0, int32_t *src1,
120- void *stream);78+ void *stream);
121-template void LaunchTDiv<int16_t, 64, 64, 64, 64, 64, 64, 64, 64, true>(int16_t *out, int16_t *src0, int16_t *src1,79+template void LaunchTDiv<int16_t, 64, 64, 64, 64, 64, 64, 64, 64, false>(int16_t *out, int16_t *src0, int16_t *src1,
122- void *stream);80+ void *stream);
123-template void LaunchTDivHalf<16, 256, 16, 256, 16, 256, 16, 256, true>(aclFloat16 *out, aclFloat16 *src0,81+template void LaunchTDivHalf<16, 256, 16, 256, 16, 256, 16, 256, false>(aclFloat16 *out, aclFloat16 *src0,
124- aclFloat16 *src1, void *stream);82+ aclFloat16 *src1, void *stream);
125template void LaunchTDivHalf<16, 64, 16, 128, 16, 128, 16, 64, false>(aclFloat16 *out, aclFloat16 *src0,83template void LaunchTDivHalf<16, 64, 16, 128, 16, 128, 16, 64, false>(aclFloat16 *out, aclFloat16 *src0,
126 aclFloat16 *src1, void *stream);84 aclFloat16 *src1, void *stream);
127template void LaunchTDiv<float, 16, 32, 16, 64, 16, 32, 16, 32, false>(float *out, float *src0, float *src1,85template void LaunchTDiv<float, 16, 32, 16, 64, 16, 32, 16, 32, false>(float *out, float *src0, float *src1,
@@ -138,3 +96,6 @@ template void LaunchTDiv<int16_t, 32, 128, 32, 128, 32, 256, 32, 127, false>(int
138 void *stream);96 void *stream);
139template void LaunchTDiv<int32_t, 16, 32, 16, 64, 16, 32, 16, 31, false>(int32_t *out, int32_t *src0, int32_t *src1,97template void LaunchTDiv<int32_t, 16, 32, 16, 64, 16, 32, 16, 31, false>(int32_t *out, int32_t *src0, int32_t *src1,
140 void *stream);98 void *stream);
99+template void LaunchTDiv<float, 2, 16, 2, 16, 2, 16, 2, 16, true>(float *out, float *src0, float *src1, void *stream);
100+template void LaunchTDivHalf<2, 32, 2, 32, 2, 32, 2, 32, true>(aclFloat16 *out, aclFloat16 *src0, aclFloat16 *src1,
101+ void *stream);
Mtests/npu/a5/src/st/testcase/tdivs/gen_data.py+18-12
@@ -14,26 +14,29 @@ import os
14import struct14import struct
15import ctypes15import ctypes
16import numpy as np16import numpy as np
17-np.random.seed(23)
18 17 
19 18 
20def gen_golden_data(param):19def gen_golden_data(param):
21- data_type = param.data_type20+ dtype = param.data_type
22 rows = param.row21 rows = param.row
23 cols = param.col22 cols = param.col
24 dst_tile_row = param.dst_tile_row23 dst_tile_row = param.dst_tile_row
25 dst_tile_col = param.dst_tile_col24 dst_tile_col = param.dst_tile_col
26 25 
27- input_arr = np.random.uniform(low=-8, high=8, size=(rows, cols)).astype(data_type)26+ if dtype in (np.int8, np.uint8, np.int16, np.uint16, np.int32, np.uint32):
28- divider = np.random.uniform(low=-8, high=8, size=(1, 1)).astype(data_type)27+ dtype_info = np.iinfo(dtype)
29- output_arr = np.zeros((dst_tile_row, dst_tile_col), dtype=data_type)28+ input_arr = np.random.randint(dtype_info.min, dtype_info.max, size=[rows, cols]).astype(dtype)
30- for i in range(rows):29+ divider = np.random.randint(dtype_info.min, dtype_info.max, size=[1, 1]).astype(dtype)
31- for j in range(cols):30+ else:
32- output_arr[i, j] = input_arr[i, j] / divider[0, 0]31+ dtype_info = np.finfo(dtype)
32+ input_arr = np.random.uniform(low=dtype_info.min, high=dtype_info.max, size=[rows, cols]).astype(dtype)
33+ divider = np.random.uniform(low=dtype_info.min, high=dtype_info.max, size=[1, 1]).astype(dtype)
34+
35+ output_arr = np.zeros((dst_tile_row, dst_tile_col), dtype=dtype)
36+ output_arr[0:rows, 0:cols] = input_arr[0:rows, 0:cols] / divider[0, 0]
33 37 
34 input_arr.tofile('input.bin')38 input_arr.tofile('input.bin')
35- with open("divider.bin", 'wb') as f:39+ divider.tofile('divider.bin')
36- f.write(struct.pack('f', np.float32(divider[0, 0])))
37 output_arr.tofile('golden.bin')40 output_arr.tofile('golden.bin')
38 41 
39 42 
@@ -46,6 +49,7 @@ class TDivsParams:
46 self.row = row49 self.row = row
47 self.col = col50 self.col = col
48 51 
52+ 
49if __name__ == "__main__":53if __name__ == "__main__":
50 case_params_list = [54 case_params_list = [
51 TDivsParams("TDIVSTest.case1", np.float32, 32, 128, 32, 64),55 TDivsParams("TDIVSTest.case1", np.float32, 32, 128, 32, 64),
@@ -53,7 +57,9 @@ if __name__ == "__main__":
53 TDivsParams("TDIVSTest.case3", np.int32, 31, 256, 31, 128),57 TDivsParams("TDIVSTest.case3", np.int32, 31, 256, 31, 128),
54 TDivsParams("TDIVSTest.case4", np.int16, 15, 192, 15, 64 * 3),58 TDivsParams("TDIVSTest.case4", np.int16, 15, 192, 15, 64 * 3),
55 TDivsParams("TDIVSTest.case5", np.float32, 7, 512, 7, 64 * 7),59 TDivsParams("TDIVSTest.case5", np.float32, 7, 512, 7, 64 * 7),
56- TDivsParams("TDIVSTest.case6", np.float32, 256, 32, 256, 16)60+ TDivsParams("TDIVSTest.case6", np.float32, 256, 32, 256, 16),
61+ TDivsParams("TDIVSTest.caseHP1", np.float32, 2, 16, 2, 16),
62+ TDivsParams("TDIVSTest.caseHP2", np.float16, 2, 32, 2, 32)
57 ]63 ]
58 64 
59 for _, case in enumerate(case_params_list):65 for _, case in enumerate(case_params_list):
@@ -62,4 +68,4 @@ if __name__ == "__main__":
62 original_dir = os.getcwd()68 original_dir = os.getcwd()
63 os.chdir(case.name)69 os.chdir(case.name)
64 gen_golden_data(case)70 gen_golden_data(case)
65- os.chdir(original_dir)71+ os.chdir(original_dir)
Mtests/npu/a5/src/st/testcase/tdivs/main.cpp+38-28
@@ -15,8 +15,13 @@ See LICENSE in the root of the software repository for the full text of the Lice
15using namespace std;15using namespace std;
16using namespace PtoTestCommon;16using namespace PtoTestCommon;
17 17 
18-template <uint32_t caseId>18+template <typename T, int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int validRow, int validCol,
19-void launchTDIVSTestCase(void *out, void *src, float scalar, aclrtStream stream);19+ bool highPrecision = false>
20+void LaunchTDivS(T *out, T *src, T scalar, void *stream);
21+ 
22+template <int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int validRow, int validCol,
23+ bool highPrecision = false>
24+void LaunchTDivSHalf(aclFloat16 *out, aclFloat16 *src, aclFloat16 scalar, void *stream);
20 25 
21class TDIVSTest : public testing::Test {26class TDIVSTest : public testing::Test {
22public:27public:
@@ -37,8 +42,9 @@ std::string GetGoldenDir()
37 return fullPath;42 return fullPath;
38}43}
39 44 
40-template <uint32_t caseId, typename T, int dstTileRow, int dstTileCol, int row, int vaildRow, int col, int srcVaildCol>45+template <typename T, int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int vaildRow, int vaildCol,
41-bool TDivSTestFramework()46+ bool isHalf = false, bool highPrecision = false>
47+void TDivSTestFramework()
42{48{
43 aclInit(nullptr);49 aclInit(nullptr);
44 aclrtSetDevice(0);50 aclrtSetDevice(0);
@@ -47,12 +53,13 @@ bool TDivSTestFramework()
47 aclrtCreateStream(&stream);53 aclrtCreateStream(&stream);
48 54 
49 size_t dstByteSize = dstTileRow * dstTileCol * sizeof(T);55 size_t dstByteSize = dstTileRow * dstTileCol * sizeof(T);
50- size_t srcByteSize = row * col * sizeof(T);56+ size_t srcByteSize = srcTileRow * srcTileCol * sizeof(T);
57+ size_t scalarByteSize = sizeof(T);
51 T *dstHost;58 T *dstHost;
52 T *srcHost;59 T *srcHost;
53 T *dstDevice;60 T *dstDevice;
54 T *srcDevice;61 T *srcDevice;
55- float scalar;62+ T scalar;
56 63 
57 aclrtMallocHost((void **)(&dstHost), dstByteSize);64 aclrtMallocHost((void **)(&dstHost), dstByteSize);
58 aclrtMallocHost((void **)(&srcHost), srcByteSize);65 aclrtMallocHost((void **)(&srcHost), srcByteSize);
@@ -61,13 +68,15 @@ bool TDivSTestFramework()
61 aclrtMalloc((void **)&srcDevice, srcByteSize, ACL_MEM_MALLOC_HUGE_FIRST);68 aclrtMalloc((void **)&srcDevice, srcByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
62 69 
63 ReadFile(GetGoldenDir() + "/input.bin", srcByteSize, srcHost, srcByteSize);70 ReadFile(GetGoldenDir() + "/input.bin", srcByteSize, srcHost, srcByteSize);
64- std::string scalar_file = GetGoldenDir() + "/divider.bin";71+ ReadFile(GetGoldenDir() + "/divider.bin", scalarByteSize, (void *)&scalar, sizeof(T));
65- std::ifstream file(scalar_file, std::ios::binary);
66- 
67- file.read(reinterpret_cast<char *>(&scalar), 4);
68- file.close();
69 aclrtMemcpy(srcDevice, srcByteSize, srcHost, srcByteSize, ACL_MEMCPY_HOST_TO_DEVICE);72 aclrtMemcpy(srcDevice, srcByteSize, srcHost, srcByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
70- launchTDIVSTestCase<caseId>(dstDevice, srcDevice, scalar, stream);73+ if constexpr (isHalf) {
74+ LaunchTDivSHalf<dstTileRow, dstTileCol, srcTileRow, srcTileCol, vaildRow, vaildCol, highPrecision>(
75+ dstDevice, srcDevice, scalar, stream);
76+ } else {
77+ LaunchTDivS<T, dstTileRow, dstTileCol, srcTileRow, srcTileCol, vaildRow, vaildCol, highPrecision>(
78+ dstDevice, srcDevice, scalar, stream);
79+ }
71 aclrtSynchronizeStream(stream);80 aclrtSynchronizeStream(stream);
72 aclrtMemcpy(dstHost, dstByteSize, dstDevice, dstByteSize, ACL_MEMCPY_DEVICE_TO_HOST);81 aclrtMemcpy(dstHost, dstByteSize, dstDevice, dstByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
73 82 
@@ -88,35 +97,36 @@ bool TDivSTestFramework()
88 ReadFile(GetGoldenDir() + "/golden.bin", dstByteSize, golden.data(), dstByteSize);97 ReadFile(GetGoldenDir() + "/golden.bin", dstByteSize, golden.data(), dstByteSize);
89 ReadFile(GetGoldenDir() + "/output.bin", dstByteSize, devFinal.data(), dstByteSize);98 ReadFile(GetGoldenDir() + "/output.bin", dstByteSize, devFinal.data(), dstByteSize);
90 99 
91- return ResultCmp<T>(golden, devFinal, 0.001f);100+ constexpr auto resPrecision = highPrecision ? 0.0000001f : 0.001f;
101+ bool ret = ResultCmp<T>(golden, devFinal, resPrecision);
102+ EXPECT_TRUE(ret);
92}103}
93 104 
94TEST_F(TDIVSTest, case1)105TEST_F(TDIVSTest, case1)
95{106{
96- bool ret = TDivSTestFramework<1, float, 32, 128, 32, 32, 64, 64>();107+ TDivSTestFramework<float, 32, 128, 32, 64, 32, 64>();
97- EXPECT_TRUE(ret);
98}108}
99- 
100TEST_F(TDIVSTest, case2)109TEST_F(TDIVSTest, case2)
101{110{
102- bool ret = TDivSTestFramework<2, aclFloat16, 63, 128, 63, 63, 64, 64>();111+ TDivSTestFramework<aclFloat16, 63, 128, 63, 64, 63, 64, true>();
103- EXPECT_TRUE(ret);
104}112}
105- 
106TEST_F(TDIVSTest, case4)113TEST_F(TDIVSTest, case4)
107{114{
108- bool ret = TDivSTestFramework<4, int16_t, 15, 192, 15, 15, 192, 192>();115+ TDivSTestFramework<int16_t, 15, 192, 15, 192, 15, 192>();
109- EXPECT_TRUE(ret);
110}116}
111- 
112TEST_F(TDIVSTest, case5)117TEST_F(TDIVSTest, case5)
113{118{
114- bool ret = TDivSTestFramework<5, float, 7, 512, 7, 7, 448, 448>();119+ TDivSTestFramework<float, 7, 512, 7, 448, 7, 448>();
115- EXPECT_TRUE(ret);
116}120}
117- 
118TEST_F(TDIVSTest, case6)121TEST_F(TDIVSTest, case6)
119{122{
120- bool ret = TDivSTestFramework<6, float, 256, 32, 256, 256, 16, 16>();123+ TDivSTestFramework<float, 256, 32, 256, 16, 256, 16>();
121- EXPECT_TRUE(ret);124+}
122-}125+TEST_F(TDIVSTest, caseHP1)
126+{
127+ TDivSTestFramework<float, 2, 16, 2, 16, 2, 16, false, true>();
128+}
129+TEST_F(TDIVSTest, caseHP2)
130+{
131+ TDivSTestFramework<aclFloat16, 2, 32, 2, 32, 2, 32, true, true>();
132+}
Mtests/npu/a5/src/st/testcase/tdivs/tdivs_kernel.cpp+27-63
@@ -15,15 +15,16 @@ See LICENSE in the root of the software repository for the full text of the Lice
15using namespace std;15using namespace std;
16using namespace pto;16using namespace pto;
17 17 
18-template <typename T, int dstTileRow, int dstTileCol, int row, int validRow, int col, int validCol>18+template <typename T, int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int validRow, int validCol,
19-PTO_INTERNAL void runTDivS(__gm__ T *out, __gm__ T *src, T scalar)19+ bool highPrecision = false>
20+__global__ AICORE void runTDIVS(__gm__ T *out, __gm__ T *src, T scalar)
20{21{
21 using DynDim2Shape = Shape<1, 1, 1, -1, -1>;22 using DynDim2Shape = Shape<1, 1, 1, -1, -1>;
22 using DynDim2Stride = pto::Stride<1, 1, -1, -1, 1>;23 using DynDim2Stride = pto::Stride<1, 1, -1, -1, 1>;
23 using GlobalData = GlobalTensor<T, DynDim2Shape, DynDim2Stride>;24 using GlobalData = GlobalTensor<T, DynDim2Shape, DynDim2Stride>;
24- GlobalData srcGlobal(src, DynDim2Shape(validRow, validCol), DynDim2Stride(row, col));25+ GlobalData srcGlobal(src, DynDim2Shape(validRow, validCol), DynDim2Stride(srcTileRow, srcTileCol));
25 GlobalData dstGlobal(out, DynDim2Shape(validRow, validCol), DynDim2Stride(dstTileRow, dstTileCol));26 GlobalData dstGlobal(out, DynDim2Shape(validRow, validCol), DynDim2Stride(dstTileRow, dstTileCol));
26- using srcTileData = Tile<TileType::Vec, T, row, col, BLayout::RowMajor, -1, -1>;27+ using srcTileData = Tile<TileType::Vec, T, srcTileRow, srcTileCol, BLayout::RowMajor, -1, -1>;
27 using dstTileData = Tile<TileType::Vec, T, dstTileRow, dstTileCol, BLayout::RowMajor, -1, -1>;28 using dstTileData = Tile<TileType::Vec, T, dstTileRow, dstTileCol, BLayout::RowMajor, -1, -1>;
28 srcTileData srcTile(validRow, validCol);29 srcTileData srcTile(validRow, validCol);
29 dstTileData dstTile(validRow, validCol);30 dstTileData dstTile(validRow, validCol);
@@ -34,7 +35,8 @@ PTO_INTERNAL void runTDivS(__gm__ T *out, __gm__ T *src, T scalar)
34 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);35 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
35 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);36 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
36#endif37#endif
37- TDIVS(dstTile, srcTile, scalar);38+ constexpr auto precisionType = highPrecision ? DivAlgorithm::HIGH_PRECISION : DivAlgorithm::DEFAULT;
39+ TDIVS<precisionType>(dstTile, srcTile, scalar);
38#ifndef __PTO_AUTO__40#ifndef __PTO_AUTO__
39 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);41 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
40 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);42 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
@@ -43,67 +45,29 @@ PTO_INTERNAL void runTDivS(__gm__ T *out, __gm__ T *src, T scalar)
43 out = dstGlobal.data();45 out = dstGlobal.data();
44}46}
45 47 
46-extern "C" __global__ AICORE void launchTDIVSCase1(__gm__ float *out, __gm__ float *src, float scalar)48+template <typename T, int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int validRow, int validCol,
49+ bool highPrecision = false>
50+void LaunchTDivS(T *out, T *src, T scalar, void *stream)
47{51{
48- runTDivS<float, 32, 128, 32, 32, 64, 64>(out, src, scalar);52+ runTDIVS<T, dstTileRow, dstTileCol, srcTileRow, srcTileCol, validRow, validCol, highPrecision>
49-}53+ <<<1, nullptr, stream>>>(out, src, scalar);
50-extern "C" __global__ AICORE void launchTDIVSCase2(__gm__ aclFloat16 *out, __gm__ aclFloat16 *src, float scalar)
51-{
52- runTDivS<half, 63, 128, 63, 63, 64, 64>((__gm__ half *)out, (__gm__ half *)src, (half)scalar);
53-}
54-extern "C" __global__ AICORE void launchTDIVSCase3(__gm__ int32_t *out, __gm__ int32_t *src, int32_t scalar)
55-{
56- runTDivS<int32_t, 31, 256, 31, 31, 128, 128>(out, src, scalar);
57-}
58-extern "C" __global__ AICORE void launchTDIVSCase4(__gm__ int16_t *out, __gm__ int16_t *src, int16_t scalar)
59-{
60- runTDivS<int16_t, 15, 192, 15, 15, 192, 192>(out, src, scalar);
61-}
62-extern "C" __global__ AICORE void launchTDIVSCase5(__gm__ float *out, __gm__ float *src, float scalar)
63-{
64- runTDivS<float, 7, 512, 7, 7, 448, 448>(out, src, scalar);
65-}
66-extern "C" __global__ AICORE void launchTDIVSCase6(__gm__ float *out, __gm__ float *src, float scalar)
67-{
68- runTDivS<float, 256, 32, 256, 256, 16, 16>(out, src, scalar);
69}54}
70 55 
71-template <uint32_t caseId>56+template <int dstTileRow, int dstTileCol, int srcTileRow, int srcTileCol, int validRow, int validCol,
72-void launchTDIVSTestCase(void *out, void *src, float scalar, aclrtStream stream)57+ bool highPrecision = false>
58+void LaunchTDivSHalf(aclFloat16 *out, aclFloat16 *src, aclFloat16 scalar, void *stream)
73{59{
74- switch (caseId) {60+ runTDIVS<half, dstTileRow, dstTileCol, srcTileRow, srcTileCol, validRow, validCol, highPrecision>
75- case 1: {61+ <<<1, nullptr, stream>>>((half *)out, (half *)src, *(half *)&scalar);
76- launchTDIVSCase1<<<1, nullptr, stream>>>((float *)out, (float *)src, scalar);
77- break;
78- }
79- case 2: {
80- launchTDIVSCase2<<<1, nullptr, stream>>>((aclFloat16 *)out, (aclFloat16 *)src, scalar);
81- break;
82- }
83- case 3: {
84- launchTDIVSCase3<<<1, nullptr, stream>>>((int32_t *)out, (int32_t *)src, scalar);
85- break;
86- }
87- case 4: {
88- launchTDIVSCase4<<<1, nullptr, stream>>>((int16_t *)out, (int16_t *)src, scalar);
89- break;
90- }
91- case 5: {
92- launchTDIVSCase5<<<1, nullptr, stream>>>((float *)out, (float *)src, scalar);
93- break;
94- }
95- case 6: {
96- launchTDIVSCase6<<<1, nullptr, stream>>>((float *)out, (float *)src, scalar);
97- break;
98- }
99- default: {
100- }
101- }
102}62}
103 63 
104-template void launchTDIVSTestCase<1>(void *out, void *src, float scalar, aclrtStream stream);64+template void LaunchTDivS<float, 32, 128, 32, 64, 32, 64>(float *out, float *src, float scalar, void *stream);
105-template void launchTDIVSTestCase<2>(void *out, void *src, float scalar, aclrtStream stream);65+template void LaunchTDivSHalf<63, 128, 63, 64, 63, 64>(aclFloat16 *out, aclFloat16 *src, aclFloat16 scalar,
106-template void launchTDIVSTestCase<3>(void *out, void *src, float scalar, aclrtStream stream);66+ void *stream);
107-template void launchTDIVSTestCase<4>(void *out, void *src, float scalar, aclrtStream stream);67+template void LaunchTDivS<int32_t, 31, 256, 31, 128, 31, 128>(int32_t *out, int32_t *src, int32_t scalar, void *stream);
108-template void launchTDIVSTestCase<5>(void *out, void *src, float scalar, aclrtStream stream);68+template void LaunchTDivS<int16_t, 15, 192, 15, 192, 15, 192>(int16_t *out, int16_t *src, int16_t scalar, void *stream);
109-template void launchTDIVSTestCase<6>(void *out, void *src, float scalar, aclrtStream stream);69+template void LaunchTDivS<float, 7, 512, 7, 448, 7, 448>(float *out, float *src, float scalar, void *stream);
70+template void LaunchTDivS<float, 256, 32, 256, 16, 256, 16>(float *out, float *src, float scalar, void *stream);
71+template void LaunchTDivS<float, 2, 16, 2, 16, 2, 16, true>(float *out, float *src, float scalar, void *stream);
72+template void LaunchTDivSHalf<2, 32, 2, 32, 2, 32, true>(aclFloat16 *out, aclFloat16 *src, aclFloat16 scalar,
73+ void *stream);
Mtests/npu/a5/src/st/testcase/trecip/gen_data.py+13-3
@@ -31,7 +31,8 @@ def gen_golden_data(case_name, param):
31 31 
32 32 
33class tunaryParams:33class tunaryParams:
34- def __init__(self, dtype, global_row, global_col, tile_row, tile_col, valid_row, valid_col, in_place = False):34+ def __init__(self, dtype, global_row, global_col, tile_row, tile_col, valid_row, valid_col, in_place=False,
35+ high_precision=False):
35 self.dtype = dtype36 self.dtype = dtype
36 self.global_row = global_row37 self.global_row = global_row
37 self.global_col = global_col38 self.global_col = global_col
@@ -40,6 +41,7 @@ class tunaryParams:
40 self.valid_row = valid_row41 self.valid_row = valid_row
41 self.valid_col = valid_col42 self.valid_col = valid_col
42 self.in_place = in_place43 self.in_place = in_place
44+ self.high_precision = high_precision
43 45 
44def generate_case_name(param):46def generate_case_name(param):
45 dtype_str = {47 dtype_str = {
@@ -49,7 +51,13 @@ def generate_case_name(param):
49 np.int32: 'int32',51 np.int32: 'int32',
50 np.int16: 'int16'52 np.int16: 'int16'
51 }[param.dtype]53 }[param.dtype]
52- return f"TRECIPTest.case_{dtype_str}_{param.global_row}x{param.global_col}_{param.tile_row}x{param.tile_col}_{param.valid_row}x{param.valid_col}_inPlace_{param.in_place}"54+ if param.high_precision:
55+ dtype_str += '_hp'
56+ inplace_flag = ''
57+ if param.in_place:
58+ inplace_flag = '_inPlace'
59+ return f"TRECIPTest.case_{dtype_str}_{param.global_row}x{param.global_col}_{param.tile_row}x{param.tile_col}_"\
60+ f"{param.valid_row}x{param.valid_col}{inplace_flag}"
53 61 
54if __name__ == "__main__":62if __name__ == "__main__":
55 # Get the absolute path of the script63 # Get the absolute path of the script
@@ -67,6 +75,8 @@ if __name__ == "__main__":
67 tunaryParams(np.float16, 64, 64, 64, 64, 64, 64, False),75 tunaryParams(np.float16, 64, 64, 64, 64, 64, 64, False),
68 tunaryParams(np.float32, 64, 64, 66, 72, 64, 64, False),76 tunaryParams(np.float32, 64, 64, 66, 72, 64, 64, False),
69 tunaryParams(np.float32, 58, 70, 66, 72, 58, 70, False),77 tunaryParams(np.float32, 58, 70, 66, 72, 58, 70, False),
78+ tunaryParams(np.float32, 2, 16, 2, 16, 2, 16, False, True),
79+ tunaryParams(np.float16, 2, 32, 2, 32, 2, 32, False, True),
70 ]80 ]
71 81 
72 for i, param in enumerate(case_params_list):82 for i, param in enumerate(case_params_list):
@@ -76,4 +86,4 @@ if __name__ == "__main__":
76 original_dir = os.getcwd()86 original_dir = os.getcwd()
77 os.chdir(case_name)87 os.chdir(case_name)
78 gen_golden_data(case_name, param)88 gen_golden_data(case_name, param)
79- os.chdir(original_dir)89+ os.chdir(original_dir)
Mtests/npu/a5/src/st/testcase/trecip/main.cpp+19-9
@@ -32,10 +32,12 @@ std::string GetGoldenDir()
32 return fullPath;32 return fullPath;
33}33}
34 34 
35-template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false>35+template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false,
36+ bool highPrecision = false>
36void LaunchTRecip(T *out, T *src, void *stream);37void LaunchTRecip(T *out, T *src, void *stream);
37 38 
38-template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false>39+template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false,
40+ bool highPrecision = false>
39void test_trecip()41void test_trecip()
40{42{
41 size_t fileSize = kGRows_ * kGCols_ * sizeof(T);43 size_t fileSize = kGRows_ * kGCols_ * sizeof(T);
@@ -57,7 +59,7 @@ void test_trecip()
57 ReadFile(GetGoldenDir() + "/input1.bin", fileSize, srcHost, fileSize);59 ReadFile(GetGoldenDir() + "/input1.bin", fileSize, srcHost, fileSize);
58 60 
59 aclrtMemcpy(srcDevice, fileSize, srcHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE);61 aclrtMemcpy(srcDevice, fileSize, srcHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE);
60- LaunchTRecip<T, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace>(dstDevice, srcDevice, stream);62+ LaunchTRecip<T, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace, highPrecision>(dstDevice, srcDevice, stream);
61 63 
62 aclrtSynchronizeStream(stream);64 aclrtSynchronizeStream(stream);
63 aclrtMemcpy(dstHost, fileSize, dstDevice, fileSize, ACL_MEMCPY_DEVICE_TO_HOST);65 aclrtMemcpy(dstHost, fileSize, dstDevice, fileSize, ACL_MEMCPY_DEVICE_TO_HOST);
@@ -89,27 +91,35 @@ void test_trecip()
89 EXPECT_TRUE(ret);91 EXPECT_TRUE(ret);
90}92}
91 93 
92-TEST_F(TRECIPTest, case_float_64x64_64x64_64x64_inPlace_True)94+TEST_F(TRECIPTest, case_float_64x64_64x64_64x64_inPlace)
93{95{
94 test_trecip<float, 64, 64, 64, 64, true>();96 test_trecip<float, 64, 64, 64, 64, true>();
95}97}
96-TEST_F(TRECIPTest, case_float_64x64_64x64_64x64_inPlace_False)98+TEST_F(TRECIPTest, case_float_64x64_64x64_64x64)
97{99{
98 test_trecip<float, 64, 64, 64, 64, false>();100 test_trecip<float, 64, 64, 64, 64, false>();
99}101}
100-TEST_F(TRECIPTest, case_half_64x64_64x64_64x64_inPlace_True)102+TEST_F(TRECIPTest, case_half_64x64_64x64_64x64_inPlace)
101{103{
102 test_trecip<aclFloat16, 64, 64, 64, 64, true>();104 test_trecip<aclFloat16, 64, 64, 64, 64, true>();
103}105}
104-TEST_F(TRECIPTest, case_half_64x64_64x64_64x64_inPlace_False)106+TEST_F(TRECIPTest, case_half_64x64_64x64_64x64)
105{107{
106 test_trecip<aclFloat16, 64, 64, 64, 64, false>();108 test_trecip<aclFloat16, 64, 64, 64, 64, false>();
107}109}
108-TEST_F(TRECIPTest, case_float_64x64_66x72_64x64_inPlace_False)110+TEST_F(TRECIPTest, case_float_64x64_66x72_64x64)
109{111{
110 test_trecip<float, 64, 64, 66, 72, false>();112 test_trecip<float, 64, 64, 66, 72, false>();
111}113}
112-TEST_F(TRECIPTest, case_float_58x70_66x72_58x70_inPlace_False)114+TEST_F(TRECIPTest, case_float_58x70_66x72_58x70)
113{115{
114 test_trecip<float, 58, 70, 66, 72, false>();116 test_trecip<float, 58, 70, 66, 72, false>();
115}117}
118+TEST_F(TRECIPTest, case_float_hp_2x16_2x16_2x16)
119+{
120+ test_trecip<float, 2, 16, 2, 16, false, true>();
121+}
122+TEST_F(TRECIPTest, case_half_hp_2x32_2x32_2x32)
123+{
124+ test_trecip<aclFloat16, 2, 32, 2, 32, false, true>();
125+}
Mtests/npu/a5/src/st/testcase/trecip/trecip_kernel.cpp+17-11
@@ -14,7 +14,8 @@ See LICENSE in the root of the software repository for the full text of the Lice
14 14 
15using namespace pto;15using namespace pto;
16 16 
17-template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false>17+template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false,
18+ bool highPrecision = false>
18__global__ AICORE void runTRecip(__gm__ T __out__ *out, __gm__ T __in__ *src)19__global__ AICORE void runTRecip(__gm__ T __out__ *out, __gm__ T __in__ *src)
19{20{
20 using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>;21 using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>;
@@ -38,7 +39,8 @@ __global__ AICORE void runTRecip(__gm__ T __out__ *out, __gm__ T __in__ *src)
38 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);39 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
39 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);40 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
40#endif41#endif
41- TRECIP(dstTile, srcTile);42+ constexpr auto precisionType = highPrecision ? DivAlgorithm::HIGH_PRECISION : DivAlgorithm::DEFAULT;
43+ TRECIP<precisionType>(dstTile, srcTile);
42#ifndef __PTO_AUTO__44#ifndef __PTO_AUTO__
43 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);45 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
44 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);46 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
@@ -47,19 +49,23 @@ __global__ AICORE void runTRecip(__gm__ T __out__ *out, __gm__ T __in__ *src)
47 out = dstGlobal.data();49 out = dstGlobal.data();
48}50}
49 51 
50-template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false>52+template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_, bool isInPlace = false,
53+ bool highPrecision = false>
51void LaunchTRecip(T *out, T *src, void *stream)54void LaunchTRecip(T *out, T *src, void *stream)
52{55{
53- if constexpr (std::is_same_v<T, aclFloat16>)56+ if constexpr (std::is_same_v<T, aclFloat16>) {
54- runTRecip<half, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace>57+ runTRecip<half, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace, highPrecision>
55 <<<1, nullptr, stream>>>((half *)(out), (half *)(src));58 <<<1, nullptr, stream>>>((half *)(out), (half *)(src));
56- else59+ } else {
57- runTRecip<T, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace><<<1, nullptr, stream>>>(out, src);60+ runTRecip<T, kGRows_, kGCols_, kTRows_, kTCols_, isInPlace, highPrecision><<<1, nullptr, stream>>>(out, src);
61+ }
58}62}
59 63 
60template void LaunchTRecip<float, 64, 64, 64, 64, true>(float *out, float *src, void *stream);64template void LaunchTRecip<float, 64, 64, 64, 64, true>(float *out, float *src, void *stream);
61-template void LaunchTRecip<float, 64, 64, 64, 64, false>(float *out, float *src, void *stream);65+template void LaunchTRecip<float, 64, 64, 64, 64>(float *out, float *src, void *stream);
62template void LaunchTRecip<aclFloat16, 64, 64, 64, 64, true>(aclFloat16 *out, aclFloat16 *src, void *stream);66template void LaunchTRecip<aclFloat16, 64, 64, 64, 64, true>(aclFloat16 *out, aclFloat16 *src, void *stream);
63-template void LaunchTRecip<aclFloat16, 64, 64, 64, 64, false>(aclFloat16 *out, aclFloat16 *src, void *stream);67+template void LaunchTRecip<aclFloat16, 64, 64, 64, 64>(aclFloat16 *out, aclFloat16 *src, void *stream);
64-template void LaunchTRecip<float, 64, 64, 66, 72, false>(float *out, float *src, void *stream);68+template void LaunchTRecip<float, 64, 64, 66, 72>(float *out, float *src, void *stream);
65-template void LaunchTRecip<float, 58, 70, 66, 72, false>(float *out, float *src, void *stream);69+template void LaunchTRecip<float, 58, 70, 66, 72>(float *out, float *src, void *stream);
70+template void LaunchTRecip<float, 2, 16, 2, 16, false, true>(float *out, float *src, void *stream);
71+template void LaunchTRecip<aclFloat16, 2, 32, 2, 32, false, true>(aclFloat16 *out, aclFloat16 *src, void *stream);
Mtests/run_st.sh+1-1
@@ -148,7 +148,7 @@ if [ "$ENABLE_A3" = "true" ]; then # A2A3
148 python3 tests/script/run_st.py $ARGS -w -v a3 -t texp -g TEXPTest.case_float_64x64_64x64_64x64_inPlace_False148 python3 tests/script/run_st.py $ARGS -w -v a3 -t texp -g TEXPTest.case_float_64x64_64x64_64x64_inPlace_False
149 python3 tests/script/run_st.py $ARGS -w -v a3 -t tabs -g TABSTest.case_float_64x64_64x64_64x64_inPlace_False149 python3 tests/script/run_st.py $ARGS -w -v a3 -t tabs -g TABSTest.case_float_64x64_64x64_64x64_inPlace_False
150 python3 tests/script/run_st.py $ARGS -w -v a3 -t tlog -g TLOGTest.case_float_64x64_64x64_64x64_inPlace_False150 python3 tests/script/run_st.py $ARGS -w -v a3 -t tlog -g TLOGTest.case_float_64x64_64x64_64x64_inPlace_False
151- python3 tests/script/run_st.py $ARGS -w -v a3 -t trecip -g TRECIPTest.case_float_64x64_64x64_64x64_inPlace_False151+ python3 tests/script/run_st.py $ARGS -w -v a3 -t trecip -g TRECIPTest.case_float_64x64_64x64_64x64
152 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case1152 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case1
153 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case4153 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case4
154 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case5154 python3 tests/script/run_st.py $ARGS -w -v a3 -t tdivs -g TDIVSTest.case5