Star1280
1391
代码
代码
Issues93
Pull Requests182
Actions
项目讨论
Wiki
分析
Star1280
1391
已合并
支持950 reduce_sum kernel输入bool类型 #3014
sakuraqqz创建于 5月29日
支持950 reduce_sum kernel输入bool类型 #3014
已合并
sakuraqqz创建于 5月29日
从sakuraqqz/ops-mathreduce_sum合入到cann/ops-mathmaster
共 7 个文件变更+107-23
Mmath/reduce_sum/op_api/aclnn_reduce_sum.cpp+8-12
@@ -258,15 +258,6 @@ aclnnStatus aclnnReduceSumGetWorkspaceSize(
258 dims = uniqueExecutor.get()->AllocIntArray(appendDim, dimDum);258 dims = uniqueExecutor.get()->AllocIntArray(appendDim, dimDum);
259 }259 }
260 260 
261- // self为bool，dtype为int64处理
262- if (selfType == op::DataType::DT_BOOL && dataType == op::DataType::DT_INT64) {
263- static const size_t maxDim = 16777216; // 2^24
264- size_t reduceDims = 1;
265- for (int i = 0; i < static_cast<int>(dims->Size()); i++) {
266- reduceDims *= shape[(*dims)[i]];
267- }
268- promoteType = (reduceDims < maxDim) ? op::DataType::DT_FLOAT : op::DataType::DT_INT64;
269- }
270 if (IsNonContiguousSupport(self, promoteType, dims)) {261 if (IsNonContiguousSupport(self, promoteType, dims)) {
271 OP_LOGD("Enter NonContigous");262 OP_LOGD("Enter NonContigous");
272 auto selfContiguous = uniqueExecutor.get()->CreateView(263 auto selfContiguous = uniqueExecutor.get()->CreateView(
@@ -287,16 +278,21 @@ aclnnStatus aclnnReduceSumGetWorkspaceSize(
287 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);278 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
288 279 
289 // 将输入self的数据类型进行转换280 // 将输入self的数据类型进行转换
290- auto selfContiguousCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());281+ 
291- CHECK_RET(selfContiguousCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);282+ const aclTensor* selfContiguousCasted = selfContiguous;
283+ if (!(selfType == op::DataType::DT_BOOL && dataType == op::DataType::DT_INT64)) {
284+ selfContiguousCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
285+ CHECK_RET(selfContiguousCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
286+ }
292 287 
293 // 调用ReduceSum算子kernel,将输入self的数据类型转换成指定的数据类型288 // 调用ReduceSum算子kernel,将输入self的数据类型转换成指定的数据类型
294 const aclTensor* reduceSumOut = nullptr;289 const aclTensor* reduceSumOut = nullptr;
295- if (dataType == op::DataType::DT_BOOL) {290+ if (dataType == op::DataType::DT_BOOL && dataType != op::DataType::DT_INT64) {
296 reduceSumOut = l0op::ReduceAny(selfContiguousCasted, dims, keepDims, uniqueExecutor.get());291 reduceSumOut = l0op::ReduceAny(selfContiguousCasted, dims, keepDims, uniqueExecutor.get());
297 } else {292 } else {
298 reduceSumOut = l0op::ReduceSumOp(selfContiguousCasted, dims, keepDims, uniqueExecutor.get());293 reduceSumOut = l0op::ReduceSumOp(selfContiguousCasted, dims, keepDims, uniqueExecutor.get());
299 }294 }
295+ 
300 CHECK_RET(reduceSumOut != nullptr, ACLNN_ERR_INNER_NULLPTR);296 CHECK_RET(reduceSumOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
301 CHECK_RET(CheckShapeAndScalarSame(reduceSumOut, out), ACLNN_ERR_PARAM_INVALID);297 CHECK_RET(CheckShapeAndScalarSame(reduceSumOut, out), ACLNN_ERR_PARAM_INVALID);
302 298 
Mmath/reduce_sum/op_api/reduce_sum_op.cpp+8-2
@@ -41,7 +41,7 @@ static const std::initializer_list<op::DataType> AICORE910B_DTYPE_SUPPORT_LIST =
41 41 
42static const std::initializer_list<op::DataType> ARCH3510_DTYPE_SUPPORT_LIST = {42static const std::initializer_list<op::DataType> ARCH3510_DTYPE_SUPPORT_LIST = {
43 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_BF16,43 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_BF16,
44- op::DataType::DT_INT64};44+ op::DataType::DT_INT64, op::DataType::DT_BOOL};
45 45 
46// 根据芯片类型、dtype判断算子是否支持走aicore46// 根据芯片类型、dtype判断算子是否支持走aicore
47static bool IsAiCoreSupport(const aclTensor* self)47static bool IsAiCoreSupport(const aclTensor* self)
@@ -100,7 +100,13 @@ static const aclTensor* ReduceSumOpAiCpu(
100const aclTensor* ReduceSumOp(const aclTensor* x, const aclIntArray* axes, bool keepDim, aclOpExecutor* executor)100const aclTensor* ReduceSumOp(const aclTensor* x, const aclIntArray* axes, bool keepDim, aclOpExecutor* executor)
101{101{
102 auto axesTensor = executor->ConvertToTensor(axes, op::ToOpDataType(ACL_INT64));102 auto axesTensor = executor->ConvertToTensor(axes, op::ToOpDataType(ACL_INT64));
103- auto out = executor->AllocTensor(x->GetDataType(), op::Format::FORMAT_ND, op::Format::FORMAT_ND);103+ aclTensor* out = nullptr;
104+ 
105+ if (x->GetDataType() == op::DataType::DT_BOOL) {
106+ out = executor->AllocTensor(op::DataType::DT_INT64, op::Format::FORMAT_ND, op::Format::FORMAT_ND);
107+ } else {
108+ out = executor->AllocTensor(x->GetDataType(), op::Format::FORMAT_ND, op::Format::FORMAT_ND);
109+ }
104 110 
105 // dim为空时，默认保留所有轴111 // dim为空时，默认保留所有轴
106 bool noopWithEmptyAxes = true;112 bool noopWithEmptyAxes = true;
Mmath/reduce_sum/op_host/arch35/reduce_sum_tiling_arch35.cpp+3-0
@@ -28,6 +28,7 @@ namespace optiling {
28static constexpr int32_t SIZE8 = 8;28static constexpr int32_t SIZE8 = 8;
29static constexpr int32_t SIZE4 = 4;29static constexpr int32_t SIZE4 = 4;
30static constexpr int32_t SIZE2 = 2;30static constexpr int32_t SIZE2 = 2;
31+static constexpr int32_t SIZE1 = 1;
31static ge::graphStatus DoTiling(gert::TilingContext* context, ReduceOpInputParam& opInput, ReduceTilingKey& key)32static ge::graphStatus DoTiling(gert::TilingContext* context, ReduceOpInputParam& opInput, ReduceTilingKey& key)
32{33{
33 ge::graphStatus status = ge::GRAPH_FAILED;34 ge::graphStatus status = ge::GRAPH_FAILED;
@@ -38,6 +39,8 @@ static ge::graphStatus DoTiling(gert::TilingContext* context, ReduceOpInputParam
38 status = Tiling4ReduceOp<ReduceSum::ReduceSumDag<float, float>::OpDag>(context, opInput, key);39 status = Tiling4ReduceOp<ReduceSum::ReduceSumDag<float, float>::OpDag>(context, opInput, key);
39 } else if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE2) {40 } else if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE2) {
40 status = Tiling4ReduceOp<ReduceSum::ReduceSumDag<half, float>::OpDag>(context, opInput, key);41 status = Tiling4ReduceOp<ReduceSum::ReduceSumDag<half, float>::OpDag>(context, opInput, key);
42+ } else if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE1) {
43+ status = Tiling4ReduceOp<ReduceSum::ReduceSumBoolDag<bool, int64_t>::OpDag>(context, opInput, key);
41 }44 }
42 OP_CHECK_IF(45 OP_CHECK_IF(
43 (status == ge::GRAPH_FAILED),46 (status == ge::GRAPH_FAILED),
Mmath/reduce_sum/op_host/config/ascend950/reduce_sum_binary.json+52-0
@@ -520,6 +520,58 @@
520 "value": null520 "value": null
521 }521 }
522 ]522 ]
523+ },
524+ {
525+ "bin_filename": "ReduceSum_bool_int64",
526+ "inputs": [
527+ {
528+ "name": "x",
529+ "index": 0,
530+ "dtype": "bool",
531+ "format": "ND",
532+ "paramType": "required",
533+ "shape": [
534+ -2
535+ ],
536+ "format_match_mode": "FormatAgnostic"
537+ },
538+ {
539+ "name": "axes",
540+ "index": 1,
541+ "dtype": "int64",
542+ "format": "ND",
543+ "paramType": "required",
544+ "shape": [
545+ -2
546+ ],
547+ "format_match_mode": "FormatAgnostic"
548+ }
549+ ],
550+ "outputs": [
551+ {
552+ "name": "y",
553+ "index": 0,
554+ "dtype": "int64",
555+ "format": "ND",
556+ "paramType": "required",
557+ "shape": [
558+ -2
559+ ],
560+ "format_match_mode": "FormatAgnostic"
561+ }
562+ ],
563+ "attrs": [
564+ {
565+ "name": "keep_dims",
566+ "dtype": "bool",
567+ "value": null
568+ },
569+ {
570+ "name": "noop_with_empty_axes",
571+ "dtype": "bool",
572+ "value": null
573+ }
574+ ]
523 }575 }
524 ]576 ]
525}577}
Mmath/reduce_sum/op_host/reduce_sum_def.cpp+11-7
@@ -15,17 +15,21 @@
15#include "register/op_def_registry.h"15#include "register/op_def_registry.h"
16 16 
17namespace ops {17namespace ops {
18-static const std::vector<ge::DataType> dataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32,18+static const std::vector<ge::DataType> dataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32,
19- ge::DT_INT64, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16,19+ ge::DT_INT64, ge::DT_BOOL, ge::DT_FLOAT16, ge::DT_FLOAT,
20- ge::DT_INT32, ge::DT_INT64};20+ ge::DT_BF16, ge::DT_INT32, ge::DT_INT64, ge::DT_BOOL};
21+
22+static const std::vector<ge::DataType> yDataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32,
23+ ge::DT_INT64, ge::DT_INT64, ge::DT_FLOAT16, ge::DT_FLOAT,
24+ ge::DT_BF16, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64};
21 25 
22static const std::vector<ge::Format> format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,26static const std::vector<ge::Format> format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
23 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,27 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
24- ge::FORMAT_ND, ge::FORMAT_ND};28+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
25 29 
26static const std::vector<ge::DataType> axesDataType = {ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,30static const std::vector<ge::DataType> axesDataType = {ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
27- ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,31+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64,
28- ge::DT_INT64, ge::DT_INT64};32+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64};
29 33 
30class ReduceSum : public OpDef {34class ReduceSum : public OpDef {
31public:35public:
@@ -35,7 +39,7 @@ public:
35 39 
36 this->Input("axes").ParamType(REQUIRED).ValueDepend(OPTIONAL).DataType(axesDataType).UnknownShapeFormat(format);40 this->Input("axes").ParamType(REQUIRED).ValueDepend(OPTIONAL).DataType(axesDataType).UnknownShapeFormat(format);
37 41 
38- this->Output("y").ParamType(REQUIRED).DataType(dataType).UnknownShapeFormat(format);42+ this->Output("y").ParamType(REQUIRED).DataType(yDataType).UnknownShapeFormat(format);
39 43 
40 this->Attr("keep_dims").AttrType(OPTIONAL).Bool(false);44 this->Attr("keep_dims").AttrType(OPTIONAL).Bool(false);
41 this->Attr("noop_with_empty_axes").AttrType(OPTIONAL).Bool(true);45 this->Attr("noop_with_empty_axes").AttrType(OPTIONAL).Bool(true);
Mmath/reduce_sum/op_kernel/arch35/reduce_sum_dag.h+13-0
@@ -36,6 +36,19 @@ struct ReduceSumDag {
36 using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;36 using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
37 using OpDag = DAGSch<Outputs, void, MemCfg>;37 using OpDag = DAGSch<Outputs, void, MemCfg>;
38};38};
39+ 
40+template <typename T, typename PromteT>
41+struct ReduceSumBoolDag {
42+ using OpCopyIn0 = Bind<Vec::CopyIn<int8_t>, Placeholder::In0<int8_t>>;
43+ using Cast0 = Bind<Vec::Cast<half, int8_t, 0>, OpCopyIn0>;
44+ using Cast1 = Bind<Vec::Cast<float, half, 0>, Cast0>;
45+ using Cast2 = Bind<Vec::Cast<int64_t, float, 1>, Cast1>;
46+ using ReduceOp0 = Bind<Vec::ReduceSumOp<int64_t>, Cast2>;
47+ using OpCopyOut = Bind<Vec::CopyOut<int64_t>, Placeholder::Out0<int64_t>, ReduceOp0>;
48+ using Outputs = Elems<OpCopyOut>;
49+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
50+ using OpDag = DAGSch<Outputs, void, MemCfg>;
51+};
39} // namespace ReduceSum52} // namespace ReduceSum
40 53 
41#endif54#endif
Mmath/reduce_sum/op_kernel/reduce_sum_apt.cpp+12-2
@@ -37,8 +37,18 @@ __global__ __aicore__ void reduce_sum(GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADD
37 REGISTER_TILING_DEFAULT(ReduceOpTilingData);37 REGISTER_TILING_DEFAULT(ReduceOpTilingData);
38 GET_TILING_DATA_WITH_STRUCT(ReduceOpTilingData, tilingData, tiling);38 GET_TILING_DATA_WITH_STRUCT(ReduceOpTilingData, tilingData, tiling);
39 TPipe pipe;39 TPipe pipe;
40- using PromoteType = __reduceType::GetPromoteType<DTYPE_X>::T;40+ using PromoteType = std::conditional_t<std::is_same_v<DTYPE_X, bool>, int64_t, __reduceType::GetPromoteType<DTYPE_X>::T>;
41- using Op = ReduceSch<REDUCE_TPL_VALUE, ReduceSum::ReduceSumDag<DTYPE_X, PromoteType>::OpDag>;41+ 
42+ // 自动选择 bool / 非 bool 分支
43+ using Op = ReduceSch<
44+ REDUCE_TPL_VALUE,
45+ std::conditional_t<
46+ std::is_same_v<DTYPE_X, bool>,
47+ typename ReduceSum::ReduceSumBoolDag<DTYPE_X, PromoteType>::OpDag,
48+ typename ReduceSum::ReduceSumDag<DTYPE_X, PromoteType>::OpDag
49+ >
50+ >;
51+ 
42 Op op(&tilingData);52 Op op(&tilingData);
43 op.Init(&pipe, x, y, userWS);53 op.Init(&pipe, x, y, userWS);
44 op.Process();54 op.Process();