Star1333
2072
代码
代码
Issues785
Pull Requests671
Actions
项目讨论
Wiki
分析
Star1333
2072
已合并
fix tiling data pointer is nullptr #2746
yu_qinfei创建于 3月16日
fix tiling data pointer is nullptr #2746
已合并
yu_qinfei创建于 3月16日
从yu_qinfei/ops-nnfix_reverseSequence_tilingdata合入到cann/ops-nn9.0.0
共 4 个文件变更+16-17
Mindex/reverse_sequence/op_host/arch35/reverse_sequence_bsa_tiling.cpp+0-1
@@ -271,7 +271,6 @@ uint64_t ReverseSequenceBSATiling::GetTilingKey() const
271 OP_LOGD("ReverseSequenceBSATiling::GetTilingKey begin");271 OP_LOGD("ReverseSequenceBSATiling::GetTilingKey begin");
272 const uint64_t tilingKey = GET_TPL_TILING_KEY(TEMPLATE_MODE, inputData_.xDtypeSize, 1);272 const uint64_t tilingKey = GET_TPL_TILING_KEY(TEMPLATE_MODE, inputData_.xDtypeSize, 1);
273 OP_LOGD(context_->GetNodeName(), "tilingKey is: [%lu]", tilingKey);273 OP_LOGD(context_->GetNodeName(), "tilingKey is: [%lu]", tilingKey);
274- printf("tiling key is :%ld", tilingKey);
275 return tilingKey;274 return tilingKey;
276}275}
277 276 
Mindex/reverse_sequence/op_kernel/arch35/reverse_sequence_struct.h+11-1
@@ -16,6 +16,16 @@
16#ifndef REVERSE_SEQUENCE_STRUCT_H16#ifndef REVERSE_SEQUENCE_STRUCT_H
17#define REVERSE_SEQUENCE_STRUCT_H17#define REVERSE_SEQUENCE_STRUCT_H
18namespace ReverseSequence {18namespace ReverseSequence {
19+ 
20+class ReverseSequenceTilingData {
21+public:
22+ int64_t batchDim;
23+ int64_t seqDim;
24+ int64_t batchSize;
25+ int64_t reverseSize;
26+ int64_t perCoreHandleNums;
27+};
28+ 
19class ReverseSequenceSimtTilingData4RegBase {29class ReverseSequenceSimtTilingData4RegBase {
20public:30public:
21 int64_t batchDim;31 int64_t batchDim;
@@ -32,7 +42,7 @@ public:
32};42};
33 43 
34class ReverseSequenceBSATilingData{44class ReverseSequenceBSATilingData{
35- public:45+public:
36 int64_t bDim;46 int64_t bDim;
37 int64_t sDim;47 int64_t sDim;
38 int64_t aDim;48 int64_t aDim;
Mindex/reverse_sequence/op_kernel/arch35/reverse_sequence_tiling_key.h+4-2
@@ -39,13 +39,15 @@ ASCENDC_TPL_SEL(
39 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),39 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
40 ASCENDC_TPL_UINT_SEL(TEMPLATE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_TEMPLATE_SIMT),40 ASCENDC_TPL_UINT_SEL(TEMPLATE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_TEMPLATE_SIMT),
41 ASCENDC_TPL_UINT_SEL(DYTPE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_DTYPE_B8, TPL_MODE_DTYPE_B16, TPL_MODE_DTYPE_B32, TPL_MODE_DTYPE_B64),41 ASCENDC_TPL_UINT_SEL(DYTPE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_DTYPE_B8, TPL_MODE_DTYPE_B16, TPL_MODE_DTYPE_B32, TPL_MODE_DTYPE_B64),
42- ASCENDC_TPL_UINT_SEL(ADDR_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_ADDR_INT32, TPL_MODE_ADDR_INT64)42+ ASCENDC_TPL_UINT_SEL(ADDR_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_ADDR_INT32, TPL_MODE_ADDR_INT64),
43+ ASCENDC_TPL_TILING_STRUCT_SEL(ReverseSequenceSimtTilingData4RegBase)
43 ),44 ),
44 ASCENDC_TPL_ARGS_SEL(45 ASCENDC_TPL_ARGS_SEL(
45 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),46 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
46 ASCENDC_TPL_UINT_SEL(TEMPLATE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_TEMPLATE_BSA),47 ASCENDC_TPL_UINT_SEL(TEMPLATE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_TEMPLATE_BSA),
47 ASCENDC_TPL_UINT_SEL(DYTPE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_DTYPE_B8, TPL_MODE_DTYPE_B16, TPL_MODE_DTYPE_B32, TPL_MODE_DTYPE_B64),48 ASCENDC_TPL_UINT_SEL(DYTPE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_DTYPE_B8, TPL_MODE_DTYPE_B16, TPL_MODE_DTYPE_B32, TPL_MODE_DTYPE_B64),
48- ASCENDC_TPL_UINT_SEL(ADDR_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_ADDR_INT64)49+ ASCENDC_TPL_UINT_SEL(ADDR_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_ADDR_INT64),
50+ ASCENDC_TPL_TILING_STRUCT_SEL(ReverseSequenceBSATilingData)
49 )51 )
50);52);
51 53 
Mindex/reverse_sequence/op_kernel/reverse_sequence_apt.cpp+1-13
@@ -25,75 +25,63 @@ __global__ __aicore__ void reverse_sequence(GM_ADDR x, GM_ADDR seq_lengths, GM_A
25{25{
26 AscendC::TPipe pipeBase;26 AscendC::TPipe pipeBase;
27 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);27 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
28- REGISTER_TILING_DEFAULT(ReverseSequenceSimtTilingData4RegBase);28+ REGISTER_TILING_DEFAULT(ReverseSequenceTilingData);
29 if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT32) {29 if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT32) {
30- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT32", ReverseSequenceSimtTilingData4RegBase);
31 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);30 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
32 ReverseSequenceSimt<int8_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);31 ReverseSequenceSimt<int8_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);
33 op.Init(x, seq_lengths, y);32 op.Init(x, seq_lengths, y);
34 op.Process();33 op.Process();
35 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64) {34 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
36- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceSimtTilingData4RegBase);
37 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);35 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
38 ReverseSequenceSimt<int8_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);36 ReverseSequenceSimt<int8_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);
39 op.Init(x, seq_lengths, y);37 op.Init(x, seq_lengths, y);
40 op.Process();38 op.Process();
41 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT32) {39 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT32) {
42- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT32", ReverseSequenceSimtTilingData4RegBase);
43 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);40 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
44 ReverseSequenceSimt<int16_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);41 ReverseSequenceSimt<int16_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);
45 op.Init(x, seq_lengths, y);42 op.Init(x, seq_lengths, y);
46 op.Process();43 op.Process();
47 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64) {44 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
48- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceSimtTilingData4RegBase);
49 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);45 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
50 ReverseSequenceSimt<int16_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);46 ReverseSequenceSimt<int16_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);
51 op.Init(x, seq_lengths, y);47 op.Init(x, seq_lengths, y);
52 op.Process();48 op.Process();
53 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT32) {49 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT32) {
54- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT32", ReverseSequenceSimtTilingData4RegBase);
55 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);50 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
56 ReverseSequenceSimt<int32_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);51 ReverseSequenceSimt<int32_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);
57 op.Init(x, seq_lengths, y);52 op.Init(x, seq_lengths, y);
58 op.Process();53 op.Process();
59 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64) {54 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
60- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceSimtTilingData4RegBase);
61 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);55 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
62 ReverseSequenceSimt<int32_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);56 ReverseSequenceSimt<int32_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);
63 op.Init(x, seq_lengths, y);57 op.Init(x, seq_lengths, y);
64 op.Process();58 op.Process();
65 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT32) {59 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT32) {
66- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT32", ReverseSequenceSimtTilingData4RegBase);
67 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);60 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
68 ReverseSequenceSimt<int64_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);61 ReverseSequenceSimt<int64_t, DTYPE_SEQ_LENGTHS, uint32_t> op(&pipeBase, &tilingData);
69 op.Init(x, seq_lengths, y);62 op.Init(x, seq_lengths, y);
70 op.Process();63 op.Process();
71 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64) {64 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
72- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_SIMT && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceSimtTilingData4RegBase);
73 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);65 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceSimtTilingData4RegBase, tilingData, tiling);
74 ReverseSequenceSimt<int64_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);66 ReverseSequenceSimt<int64_t, DTYPE_SEQ_LENGTHS, uint64_t> op(&pipeBase, &tilingData);
75 op.Init(x, seq_lengths, y);67 op.Init(x, seq_lengths, y);
76 op.Process();68 op.Process();
77 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64) {69 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
78- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B64 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceBSATilingData);
79 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);70 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);
80 ReverseSequenceBSA<int64_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);71 ReverseSequenceBSA<int64_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);
81 op.Init(x, seq_lengths, y);72 op.Init(x, seq_lengths, y);
82 op.Process();73 op.Process();
83 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64) {74 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
84- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B32 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceBSATilingData);
85 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);75 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);
86 ReverseSequenceBSA<int32_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);76 ReverseSequenceBSA<int32_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);
87 op.Init(x, seq_lengths, y);77 op.Init(x, seq_lengths, y);
88 op.Process();78 op.Process();
89 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64) {79 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
90- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B16 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceBSATilingData);
91 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);80 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);
92 ReverseSequenceBSA<int16_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);81 ReverseSequenceBSA<int16_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);
93 op.Init(x, seq_lengths, y);82 op.Init(x, seq_lengths, y);
94 op.Process();83 op.Process();
95 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64) {84 } else if constexpr (TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64) {
96- REGISTER_TILING_FOR_TILINGKEY("TEMPLATE_MODE == TPL_MODE_TEMPLATE_BSA && DYTPE_MODE == TPL_MODE_DTYPE_B8 && ADDR_MODE == TPL_MODE_ADDR_INT64", ReverseSequenceBSATilingData);
97 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);85 GET_TILING_DATA_WITH_STRUCT(ReverseSequenceBSATilingData, tilingData, tiling);
98 ReverseSequenceBSA<int8_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);86 ReverseSequenceBSA<int8_t, DTYPE_SEQ_LENGTHS> op(&pipeBase, &tilingData);
99 op.Init(x, seq_lengths, y);87 op.Init(x, seq_lengths, y);