Star1309
1430
代码
代码
Issues118
Pull Requests248
Actions
项目讨论
Wiki
分析
Star1309
1430
已合并
[Reduce模板] Reduce算子模板tilingKey宏统一修改 #1536
Sunee创建于 3月7日
[Reduce模板] Reduce算子模板tilingKey宏统一修改 #1536
已合并
Sunee创建于 3月7日
从Sunee/ops-mathfeature/reduceTmpl合入到cann/ops-mathmaster
共 16 个文件变更+30-23
Mcmake/third_party/opbase.cmake+1-1
@@ -24,7 +24,7 @@ else()
24 FetchContent_Declare(24 FetchContent_Declare(
25 opbase25 opbase
26 GIT_REPOSITORY https://gitcode.com/cann/opbase.git26 GIT_REPOSITORY https://gitcode.com/cann/opbase.git
27- GIT_TAG 929902e577bb809a377a6ab17cb7b00a4dfddc1a27+ GIT_TAG 07383cb90f23b3c4e9b1e16eff40a048c888555f
28 GIT_PROGRESS TRUE28 GIT_PROGRESS TRUE
29 SOURCE_DIR ${CANN_3RD_LIB_PATH}/opbase)29 SOURCE_DIR ${CANN_3RD_LIB_PATH}/opbase)
30 30 
Mmath/bias_add_grad/op_host/arch35/bias_add_grad_tiling_arch35.cpp+2-2
@@ -91,8 +91,8 @@ ge::graphStatus Tiling4BiasAddGradForAscendC(gert::TilingContext* context)
91 OP_CHECK_IF(91 OP_CHECK_IF(
92 (DoTiling(context, compileInfo, opInput, key) == ge::GRAPH_FAILED),92 (DoTiling(context, compileInfo, opInput, key) == ge::GRAPH_FAILED),
93 OP_LOGE(context->GetNodeName(), "DoTiling Failed for BiasAddGrad"), return ge::GRAPH_FAILED);93 OP_LOGE(context->GetNodeName(), "DoTiling Failed for BiasAddGrad"), return ge::GRAPH_FAILED);
94- 94+ uint64_t tilingKey;
95- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);95+ GEN_REDUCE_TILING_KEY(tilingKey, key);
96 OP_LOGI(96 OP_LOGI(
97 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,97 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
98 key.loopARCount, key.loopInnerARCount, tilingKey);98 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/dot/op_host/arch35/dot_tiling_arch35.cpp+2-2
@@ -102,8 +102,8 @@ static ge::graphStatus GenInput(gert::TilingContext* context, ReduceOpInputParam
102 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),102 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
103 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceSum"),103 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceSum"),
104 return ge::GRAPH_FAILED);104 return ge::GRAPH_FAILED);
105- 105+ uint64_t tilingKey;
106- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);106+ GEN_REDUCE_TILING_KEY(tilingKey, key);
107 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",107 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",
108 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);108 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);
109 context->SetTilingKey(tilingKey);109 context->SetTilingKey(tilingKey);
Mmath/kl_div_v2/op_host/arch35/kl_div_v2_tiling_arch35.cpp+2-2
@@ -157,8 +157,8 @@ static ge::graphStatus Tiling4KLDivV2AscendC(gert::TilingContext *context, const
157 OP_CHECK_IF(it == STR_2_INT.end(),157 OP_CHECK_IF(it == STR_2_INT.end(),
158 OP_LOGE(context->GetNodeName(), "reduction Failed for KLDivV2"),158 OP_LOGE(context->GetNodeName(), "reduction Failed for KLDivV2"),
159 return ge::GRAPH_FAILED);159 return ge::GRAPH_FAILED);
160- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount, it->second,160+ uint64_t tilingKey;
161- static_cast<uint32_t>(logTarget));161+ GEN_REDUCE_TILING_KEY(tilingKey, key, it->second, static_cast<uint32_t>(logTarget));
162 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, reduction is:%u, logTarget is %u Tiling Key is:%lu",162 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, reduction is:%u, logTarget is %u Tiling Key is:%lu",
163 key.patternID, key.loopARCount, key.loopInnerARCount, it->second, static_cast<uint32_t>(logTarget), tilingKey);163 key.patternID, key.loopARCount, key.loopInnerARCount, it->second, static_cast<uint32_t>(logTarget), tilingKey);
164 context->SetTilingKey(tilingKey);164 context->SetTilingKey(tilingKey);
Mmath/reduce_all/op_host/arch35/reduce_all_tiling_arch35.cpp+2-1
@@ -48,7 +48,8 @@ static ge::graphStatus Tiling4ReduceAll(gert::TilingContext* context)
48 OP_CHECK_IF(48 OP_CHECK_IF(
49 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),49 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
50 OP_LOGE(context->GetNodeName(), "Tiling For ReduceAll Failed"), return ge::GRAPH_FAILED);50 OP_LOGE(context->GetNodeName(), "Tiling For ReduceAll Failed"), return ge::GRAPH_FAILED);
51- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);51+ uint64_t tilingKey;
52+ GEN_REDUCE_TILING_KEY(tilingKey, key);
52 OP_LOGI(53 OP_LOGI(
53 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,54 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
54 key.loopARCount, key.loopInnerARCount, tilingKey);55 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_any/op_host/arch35/reduce_any_tiling_arch35.cpp+2-1
@@ -55,7 +55,8 @@ static ge::graphStatus Tiling4ReduceAny(gert::TilingContext* context)
55 OP_CHECK_IF(55 OP_CHECK_IF(
56 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),56 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
57 OP_LOGE(context->GetNodeName(), "Tiling For ReduceAny Failed"), return ge::GRAPH_FAILED);57 OP_LOGE(context->GetNodeName(), "Tiling For ReduceAny Failed"), return ge::GRAPH_FAILED);
58- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);58+ uint64_t tilingKey;
59+ GEN_REDUCE_TILING_KEY(tilingKey, key);
59 OP_LOGI(60 OP_LOGI(
60 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,61 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
61 key.loopARCount, key.loopInnerARCount, tilingKey);62 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_log_sum_exp/op_host/arch35/reduce_log_sum_exp_tiling_arch35.cpp+2-1
@@ -54,7 +54,8 @@ ge::graphStatus Tiling4ReduceLogSumExp(gert::TilingContext* context)
54 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),54 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
55 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceLogSumExp"),55 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceLogSumExp"),
56 return ge::GRAPH_FAILED);56 return ge::GRAPH_FAILED);
57- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);57+ uint64_t tilingKey;
58+ GEN_REDUCE_TILING_KEY(tilingKey, key);
58 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",59 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",
59 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);60 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);
60 context->SetTilingKey(tilingKey);61 context->SetTilingKey(tilingKey);
Mmath/reduce_log_sum_exp/op_kernel/reduce_log_sum_exp_apt.cpp+1-1
@@ -36,7 +36,7 @@ __aicore__ inline constexpr T GetDumpValue()
36 }36 }
37}37}
38 38 
39-template <uint32_t PatternID, uint32_t LoopARCount, uint32_t LoopInnerARCount>39+template <REDUCE_TPL_PARAM>
40__global__ __aicore__ void reduce_log_sum_exp(GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)40__global__ __aicore__ void reduce_log_sum_exp(GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
41{41{
42 if (g_coreType == AIC) {42 if (g_coreType == AIC) {
Mmath/reduce_max/op_host/arch35/reduce_max_tiling_arch35.cpp+2-1
@@ -76,7 +76,8 @@ static ge::graphStatus Tiling4ReduceMax(gert::TilingContext* context)
76 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),76 OP_CHECK_IF((DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
77 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMax"),77 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMax"),
78 return ge::GRAPH_FAILED);78 return ge::GRAPH_FAILED);
79- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);79+ uint64_t tilingKey;
80+ GEN_REDUCE_TILING_KEY(tilingKey, key);
80 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",81 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",
81 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);82 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);
82 context->SetTilingKey(tilingKey);83 context->SetTilingKey(tilingKey);
Mmath/reduce_mean/op_host/arch35/reduce_mean_tiling_arch35.cpp+2-1
@@ -73,7 +73,8 @@ static ge::graphStatus Tiling4ReduceMean(gert::TilingContext* context)
73 OP_CHECK_IF(73 OP_CHECK_IF(
74 (DoTilingAscendC(context, compile_info, opInput, key) == ge::GRAPH_FAILED),74 (DoTilingAscendC(context, compile_info, opInput, key) == ge::GRAPH_FAILED),
75 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMean"), return ge::GRAPH_FAILED);75 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMean"), return ge::GRAPH_FAILED);
76- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);76+ uint64_t tilingKey;
77+ GEN_REDUCE_TILING_KEY(tilingKey, key);
77 OP_LOGI(78 OP_LOGI(
78 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,79 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
79 key.loopARCount, key.loopInnerARCount, tilingKey);80 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_min/op_host/arch35/reduce_min_tiling_arch35.cpp+2-1
@@ -76,7 +76,8 @@ static ge::graphStatus Tiling4ReduceMin(gert::TilingContext* context)
76 OP_CHECK_IF(76 OP_CHECK_IF(
77 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),77 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
78 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMin"), return ge::GRAPH_FAILED);78 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceMin"), return ge::GRAPH_FAILED);
79- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);79+ uint64_t tilingKey;
80+ GEN_REDUCE_TILING_KEY(tilingKey, key);
80 OP_LOGI(81 OP_LOGI(
81 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,82 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
82 key.loopARCount, key.loopInnerARCount, tilingKey);83 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_prod/op_host/arch35/reduce_prod_tiling_arch35.cpp+2-2
@@ -74,8 +74,8 @@ static ge::graphStatus Tiling4ReduceProd(gert::TilingContext* context)
74 OP_CHECK_IF(74 OP_CHECK_IF(
75 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),75 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
76 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceProd"), return ge::GRAPH_FAILED);76 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceProd"), return ge::GRAPH_FAILED);
77- 77+ uint64_t tilingKey;
78- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);78+ GEN_REDUCE_TILING_KEY(tilingKey, key);
79 OP_LOGI(79 OP_LOGI(
80 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,80 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
81 key.loopARCount, key.loopInnerARCount, tilingKey);81 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_std_v2/op_host/arch35/reduce_std_v2_tiling_arch35.cpp+2-2
@@ -29,8 +29,8 @@ static ge::graphStatus Tiling4ReduceStdV2(gert::TilingContext* context) {
29 OP_CHECK_IF((tiling.RunTiling(key) != ge::GRAPH_SUCCESS),29 OP_CHECK_IF((tiling.RunTiling(key) != ge::GRAPH_SUCCESS),
30 OP_LOGE(context->GetNodeName(), "RunTiling Failed for ReduceStdV2"),30 OP_LOGE(context->GetNodeName(), "RunTiling Failed for ReduceStdV2"),
31 return ge::GRAPH_FAILED);31 return ge::GRAPH_FAILED);
32- 32+ uint64_t tilingKey;
33- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);33+ GEN_REDUCE_TILING_KEY(tilingKey, key);
34 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",34 OP_LOGI(context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu",
35 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);35 key.patternID, key.loopARCount, key.loopInnerARCount, tilingKey);
36 context->SetTilingKey(tilingKey);36 context->SetTilingKey(tilingKey);
Mmath/reduce_sum/op_host/arch35/reduce_sum_tiling_arch35.cpp+2-2
@@ -72,8 +72,8 @@ static ge::graphStatus Tiling4ReduceSum(gert::TilingContext* context)
72 OP_CHECK_IF(72 OP_CHECK_IF(
73 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),73 (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
74 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceSum"), return ge::GRAPH_FAILED);74 OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceSum"), return ge::GRAPH_FAILED);
75- 75+ uint64_t tilingKey;
76- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);76+ GEN_REDUCE_TILING_KEY(tilingKey, key);
77 OP_LOGI(77 OP_LOGI(
78 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,78 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
79 key.loopARCount, key.loopInnerARCount, tilingKey);79 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/reduce_var/op_host/arch35/reduce_var_tiling.cpp+2-2
@@ -1140,8 +1140,8 @@ static ge::graphStatus Tiling4ReduceVar(gert::TilingContext* context)
1140 OP_CHECK_IF((tiling.RunTiling(key) != ge::GRAPH_SUCCESS),1140 OP_CHECK_IF((tiling.RunTiling(key) != ge::GRAPH_SUCCESS),
1141 OP_LOGE(context->GetNodeName(), "RunTiling Failed for ReduceVar"),1141 OP_LOGE(context->GetNodeName(), "RunTiling Failed for ReduceVar"),
1142 return ge::GRAPH_FAILED);1142 return ge::GRAPH_FAILED);
1143- 1143+ uint64_t tilingKey;
1144- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);1144+ GEN_REDUCE_TILING_KEY(tilingKey, key);
1145 OP_LOGI(1145 OP_LOGI(
1146 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,1146 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
1147 key.loopARCount, key.loopInnerARCount, tilingKey);1147 key.loopARCount, key.loopInnerARCount, tilingKey);
Mmath/square_sum_v1/op_host/arch35/square_sum_v1_tiling.cpp+2-1
@@ -65,7 +65,8 @@ static ge::graphStatus Tiling4SquareSumV1AscendC(gert::TilingContext* context)
65 OP_CHECK_IF(65 OP_CHECK_IF(
66 (DoTilingAscendC(context, opInput, key) == ge::GRAPH_FAILED),66 (DoTilingAscendC(context, opInput, key) == ge::GRAPH_FAILED),
67 OP_LOGE(context->GetNodeName(), "DoTiling Failed for SquareSumV1"), return ge::GRAPH_FAILED);67 OP_LOGE(context->GetNodeName(), "DoTiling Failed for SquareSumV1"), return ge::GRAPH_FAILED);
68- const uint64_t tilingKey = GET_TPL_TILING_KEY(key.patternID, key.loopARCount, key.loopInnerARCount);68+ uint64_t tilingKey;
69+ GEN_REDUCE_TILING_KEY(tilingKey, key);
69 OP_LOGI(70 OP_LOGI(
70 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,71 context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, Tiling Key is:%lu", key.patternID,
71 key.loopARCount, key.loopInnerARCount, tilingKey);72 key.loopARCount, key.loopInnerARCount, tilingKey);