From b1042e8d78421ab2bc73cee62b4d01faaadb5f60 Mon Sep 17 00:00:00 2001 From: 13611516375 Date: Mon, 3 Jul 2023 11:11:25 +0800 Subject: [PATCH] =?UTF-8?q?=E4=BD=BF=E7=94=A8ascend=20c=E5=BC=80=E5=8F=91c?= =?UTF-8?q?osh=E6=A0=B7=E4=BE=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../ai_core/op_host/cosh_ascendc.cpp | 115 ++++++++++++++++++ .../ai_core/op_host/cosh_ascendc_tiling.h | 44 +++++++ .../ai_core/op_kernel/cosh_ascendc.cpp | 111 +++++++++++++++++ .../ops/cosh_ascendc/framework/cosh_plugin.cc | 24 ++++ .../tests/cosh_ascendc/st/cosh_ascendc.json | 46 +++++++ .../tests/cosh_ascendc/st/test_cosh_data.py | 20 +++ .../ut/tik2/cosh_ascendc_gen_data.py | 53 ++++++++ .../cosh_ascendc/ut/tik2/cosh_ascendc_ut.cc | 63 ++++++++++ scripts/gen_ops_filter.sh | 4 +- 9 files changed, 479 insertions(+), 1 deletion(-) create mode 100644 community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc.cpp create mode 100644 community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc_tiling.h create mode 100644 community/ops/cosh_ascendc/ai_core/op_kernel/cosh_ascendc.cpp create mode 100644 community/ops/cosh_ascendc/framework/cosh_plugin.cc create mode 100644 community/tests/cosh_ascendc/st/cosh_ascendc.json create mode 100644 community/tests/cosh_ascendc/st/test_cosh_data.py create mode 100644 community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_gen_data.py create mode 100644 community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_ut.cc diff --git a/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc.cpp b/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc.cpp new file mode 100644 index 00000000..5bb62a04 --- /dev/null +++ b/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc.cpp @@ -0,0 +1,115 @@ +/** + * Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "cosh_ascendc_tiling.h" +#include "register/op_def_registry.h" + +namespace optiling { +struct TilingCompileInfo { + int64_t ub_size; +}; + +static ge::graphStatus TilingFunc(gert::TilingContext* context) +{ + CoshAscendcTilingData tiling; + context->SetBlockDim(8); + uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize(); + tiling.set_blockDim(8); + tiling.set_totalLength(totalLength); + tiling.set_tileNum(8); + tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); + context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); + return ge::GRAPH_SUCCESS; +} + +static ge::graphStatus TilingPrepare(gert::TilingParseContext* context) +{ + return ge::GRAPH_SUCCESS; +} + +static int32_t CheckOpSupport(const ge::Operator &op, ge::AscendString &result) +{ + std::string res_json_str = "{\"ret_code\": \"0\",\"reason\": \"check_supported_stub\"}"; + result = ge::AscendString(res_json_str.c_str()); + return 1; +} +} // namespace optiling + +namespace ge { +ge::graphStatus CoshAscendcInferShape(gert::InferShapeContext* context) +{ + const gert::Shape* x1_shape = context->GetInputShape(0); + gert::Shape* y_shape = context->GetOutputShape(0); + *y_shape = *x1_shape; + return GRAPH_SUCCESS; +} + +ge::graphStatus CoshAscendcInferShapeRange(gert::InferShapeRangeContext* context) +{ + const gert::Range* x1_shape_range = context->GetInputShapeRange(0); + gert::Range* y_shape_range = context->GetOutputShapeRange(0); + *y_shape_range = *x1_shape_range; + return GRAPH_SUCCESS; +} + +ge::graphStatus CoshAscendcInferDataType(gert::InferDataTypeContext* context) +{ + const ge::DataType x1_datatype = context->GetInputDataType(0); + context->SetOutputDataType(0, x1_datatype); + return GRAPH_SUCCESS; +} +} + +namespace ops { +class CoshAscendc : public OpDef { +public: + explicit CoshAscendc(const char* name) : OpDef(name) + { + this->Input("x") + .ParamType(REQUIRED) + .DataType({ge::DT_FLOAT16}) + .Format({ge::FORMAT_ND}) + .UnknownShapeFormat({ge::FORMAT_ND}); + this->Output("y") + .ParamType(REQUIRED) + .DataType({ge::DT_FLOAT16}) + .Format({ge::FORMAT_ND}) + .UnknownShapeFormat({ge::FORMAT_ND}); + + this->SetInferShape(ge::CoshAscendcInferShape) + .SetInferShapeRange(ge::CoshAscendcInferShapeRange) + .SetInferDataType(ge::CoshAscendcInferDataType); + + this->AICore() + .SetTiling(optiling::TilingFunc) + .SetTilingParse(optiling::TilingPrepare) + .SetCheckSupport(optiling::CheckOpSupport); + + OpAICoreConfig aicConfig; + aicConfig.AsyncFlag(true) + .DynamicCompileStaticFlag(true) + .DynamicFormatFlag(true) + .DynamicRankSupportFlag(true) + .DynamicShapeSupportFlag(true) + .NeedCheckSupportFlag(true) + .PrecisionReduceFlag(true) + .RangeLimitValue("limited"); + this->AICore().AddConfig("ascend910", aicConfig); + } +}; + +OP_ADD(CoshAscendc, optiling::TilingCompileInfo); +} // namespace ops diff --git a/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc_tiling.h b/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc_tiling.h new file mode 100644 index 00000000..a8f95906 --- /dev/null +++ b/community/ops/cosh_ascendc/ai_core/op_host/cosh_ascendc_tiling.h @@ -0,0 +1,44 @@ +/** + * Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef ADD_TIK2_TILING_H +#define ADD_TIK2_TILING_H + +#ifdef TIK2_UT + +struct CoshAscendcTilingData { + uint32_t blockDim; + uint32_t totalLength; + uint32_t tileNum; +}; + +#define GET_TILING_DATA(tilingData, tilingPointer) \ + CoshAscendcTilingData *tilingDataPointer = reinterpret_cast((uint8_t *)(tilingPointer)); \ + CoshAscendcTilingData tilingData(*tilingDataPointer); + +#else // TIK2_UT +#include "register/tilingdata_base.h" +namespace optiling { +BEGIN_TILING_DATA_DEF(CoshAscendcTilingData) + TILING_DATA_FIELD_DEF(uint32_t, blockDim); + TILING_DATA_FIELD_DEF(uint32_t, totalLength); + TILING_DATA_FIELD_DEF(uint32_t, tileNum); +END_TILING_DATA_DEF; + +REGISTER_TILING_DATA_CLASS(CoshAscendc, CoshAscendcTilingData) +} +#endif // TIK2_UT +#endif // ADD_TIK2_TILING_H diff --git a/community/ops/cosh_ascendc/ai_core/op_kernel/cosh_ascendc.cpp b/community/ops/cosh_ascendc/ai_core/op_kernel/cosh_ascendc.cpp new file mode 100644 index 00000000..dbc73525 --- /dev/null +++ b/community/ops/cosh_ascendc/ai_core/op_kernel/cosh_ascendc.cpp @@ -0,0 +1,111 @@ +/** + * Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel_operator.h" +#ifdef TIK2_UT +#include "../op_host/cosh_ascendc_tiling.h" +#endif +using namespace tik2; + +constexpr int32_t BUFFER_NUM = 2; + +class KernelCosh { +public: + __aicore__ inline KernelCosh() {} + __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, + uint32_t blockDim, uint32_t totalLength, uint32_t tileNum) + { + this->blockLength = totalLength / blockDim; + this->tileNum = tileNum; + this->tileLength = this->blockLength / tileNum / BUFFER_NUM; + + xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * this->blockLength, this->blockLength); + yGm.SetGlobalBuffer((__gm__ half*)y + block_idx * this->blockLength, this->blockLength); + pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half)); + pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(half)); + pipe.InitBuffer(tmpBuffer, this->tileLength * sizeof(half)); + } + __aicore__ inline void Process() + { + int32_t loopCount = this->tileNum * BUFFER_NUM; + for (int32_t i = 0; i < loopCount; i++) { + CopyIn(i); + Compute(i); + CopyOut(i); + } + } + +private: + __aicore__ inline void CopyIn(int32_t progress) + { + LocalTensor xLocal = inQueueX.AllocTensor(); + DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength); + inQueueX.EnQue(xLocal); + } + __aicore__ inline void Compute(int32_t progress) + { + // cosh(x) = (exp(x) + exp(-x)) / 2.0 + LocalTensor xLocal = inQueueX.DeQue(); + LocalTensor yLocal = outQueueY.AllocTensor(); + LocalTensor tmpTensor = tmpBuffer.Get(); + // exp(x) + Exp(yLocal, xLocal, this->tileLength); + half inputVal(-1.0); + Duplicate(tmpTensor, inputVal, this->tileLength); + Div(tmpTensor,xLocal,tmpTensor,this->tileLength); + // exp(-x) + Exp(tmpTensor, tmpTensor, this->tileLength); + // exp(x) + exp(-x) + yLocal = yLocal + tmpTensor; + half inputVal2(2.0); + Duplicate(tmpTensor, inputVal2, this->tileLength); + // (exp(x) + exp(-x)) / 2.0 + Div(yLocal,yLocal,tmpTensor,this->tileLength); + outQueueY.EnQue(yLocal); + inQueueX.FreeTensor(xLocal); + } + __aicore__ inline void CopyOut(int32_t progress) + { + LocalTensor yLocal = outQueueY.DeQue(); + DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength); + outQueueY.FreeTensor(yLocal); + } + +private: + TPipe pipe; + TQue inQueueX; + TQue outQueueY; + TBuf tmpBuffer; + GlobalTensor xGm, yGm; + uint32_t blockLength; + uint32_t tileNum; + uint32_t tileLength; +}; + +extern "C" __global__ __aicore__ void cosh_ascendc(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* tiling) +{ + GET_TILING_DATA(tilingData, tiling); + KernelCosh op; + op.Init(x, y, tilingData.blockDim, tilingData.totalLength, tilingData.tileNum); + op.Process(); +} + +#ifndef __CCE_KT_TEST__ +void cosh_ascendc(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* tiling) +{ + cosh_ascendc<<>>(x, y, tiling); +} +#endif diff --git a/community/ops/cosh_ascendc/framework/cosh_plugin.cc b/community/ops/cosh_ascendc/framework/cosh_plugin.cc new file mode 100644 index 00000000..66944dae --- /dev/null +++ b/community/ops/cosh_ascendc/framework/cosh_plugin.cc @@ -0,0 +1,24 @@ +/** + * Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "register/register.h" + +namespace domi { +REGISTER_CUSTOM_OP("CoshAscnedc") + .FrameworkType(TENSORFLOW) + .OriginOpType("Cosh") + .ParseParamsByOperatorFn(AutoMappingByOpFn); +} // namespace domi \ No newline at end of file diff --git a/community/tests/cosh_ascendc/st/cosh_ascendc.json b/community/tests/cosh_ascendc/st/cosh_ascendc.json new file mode 100644 index 00000000..3ce408b6 --- /dev/null +++ b/community/tests/cosh_ascendc/st/cosh_ascendc.json @@ -0,0 +1,46 @@ +[ + { + "case_name":"Test_CoshAscendC_001", + "op":"CoshAscendc", + "calc_expect_func_file": "test_cosh_data.py:calc_expect_func", + "input_desc":[ + { + "name": "x", + "format":[ + "ND" + ], + "type":[ + "float16" + ], + "shape":[ + 8, + 1024 + ], + "value_range":[ + [ + 0.1, + 1.0 + ] + ], + "data_distribute":[ + "uniform" + ] + } + ], + "output_desc":[ + { + "name": "y", + "format":[ + "ND" + ], + "type":[ + "float16" + ], + "shape":[ + 8, + 1024 + ] + } + ] + } +] \ No newline at end of file diff --git a/community/tests/cosh_ascendc/st/test_cosh_data.py b/community/tests/cosh_ascendc/st/test_cosh_data.py new file mode 100644 index 00000000..9680de85 --- /dev/null +++ b/community/tests/cosh_ascendc/st/test_cosh_data.py @@ -0,0 +1,20 @@ +# Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# ============================================================================ +import numpy as np + + +def calc_expect_func(x, y): + res = np.cosh(x["value"]) + return [res] \ No newline at end of file diff --git a/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_gen_data.py b/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_gen_data.py new file mode 100644 index 00000000..2dd6523e --- /dev/null +++ b/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_gen_data.py @@ -0,0 +1,53 @@ +# Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# ============================================================================ + +import numpy as np +import os + + +def write_file_txt(file_name, data, fmt="%s"): + if (file_name is None): + print("file name is none, do not write data to file") + return + dir_name = os.path.dirname(file_name) + if not os.path.exists(dir_name): + os.makedirs(dir_name) + np.savetxt(file_name, data.flatten(), fmt=fmt, delimiter='', newline='\n') + + +np.random.seed(677) + +one_repeat_calcount = 128 +block_dim_imm = 8 +tile_num_imm = 8 +double_buffer_imm = 2 +total_length_imm = block_dim_imm * \ + one_repeat_calcount * tile_num_imm * double_buffer_imm + +block_dim = np.array(block_dim_imm, dtype=np.uint32) +total_length = np.array(total_length_imm, dtype=np.uint32) +tile_num = np.array(tile_num_imm, dtype=np.uint32) +tiling = (block_dim, total_length, tile_num) +tiling_data = b''.join(x.tobytes() for x in tiling) + +input_x = np.random.uniform(0, 1, [total_length_imm,]).astype(np.float16) +golden = np.cosh(input_x) + +write_file_txt("cosh_ascendc/data/golden.txt", golden, fmt="%s") +with open('cosh_ascendc/data/tiling.bin', "wb") as f: + f.write(tiling_data) + +input_x.tofile("cosh_ascendc/data/input_x.bin") +write_file_txt("cosh_ascendc/data/input_x.txt", input_x, fmt="%s") \ No newline at end of file diff --git a/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_ut.cc b/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_ut.cc new file mode 100644 index 00000000..abc940f3 --- /dev/null +++ b/community/tests/cosh_ascendc/ut/tik2/cosh_ascendc_ut.cc @@ -0,0 +1,63 @@ +/** + * Copyright (c) 2023-2023 Huawei Technologies Co., Ltd. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include +#include +#include "tik2_ut_util.h" +#include "tikicpulib.h" + +extern "C" void cosh_ascendc(uint8_t* x, uint8_t* y, uint8_t* tiling); + +class SinhTest : public testing::Test { +protected: + static void SetUpTestCase() { + std::cout << "sinh test SetUp" << std::endl; +} + + static void TearDownTestCase() { + std::cout << "sinh test TearDown" << std::endl; + } +}; + +TEST_F(SinhTest, sinh_test_case_1) { + size_t tilingSize = 3 * sizeof(uint32_t); + uint8_t* tiling = (uint8_t*)tik2::GmAlloc(tilingSize); + ReadFile(ktestcaseFilePath + "cosh_ascendc/data/tiling.bin", tilingSize, tiling, tilingSize); + + uint32_t blockDim = (*(const uint32_t*)(tiling)); + size_t inputByteSize = blockDim * 2048 * sizeof(uint16_t); + size_t outputByteSize = blockDim * 2048 * sizeof(uint16_t); + + uint8_t* x = (uint8_t*)tik2::GmAlloc(inputByteSize); + uint8_t* y = (uint8_t*)tik2::GmAlloc(outputByteSize); + + ReadFile(ktestcaseFilePath + "cosh_ascendc/data/input_x.bin", inputByteSize, x, inputByteSize); + + + ICPU_RUN_KF(cosh_ascendc, blockDim, x, y, tiling); + + WriteFile(ktestcaseFilePath + "cosh_ascendc/data/output_y.bin", y, outputByteSize); + + size_t elementsNum = blockDim * 2048 ; + half* golden = new half[elementsNum]; + auto goldenFilePath = ktestcaseFilePath + "cosh_ascendc/data/golden.txt"; + ReadFile(goldenFilePath, golden, elementsNum); + bool compare = CompareResult((half*)y, golden, elementsNum); + tik2::GmFree((void*)x); + tik2::GmFree((void*)y); + tik2::GmFree((void*)tiling); + + EXPECT_EQ(compare, true); +} \ No newline at end of file diff --git a/scripts/gen_ops_filter.sh b/scripts/gen_ops_filter.sh index e55aa998..a211285e 100755 --- a/scripts/gen_ops_filter.sh +++ b/scripts/gen_ops_filter.sh @@ -57,7 +57,9 @@ for file in ${ini_files} ; do else isHeavy="false" fi - add_ops ${name} ${isHeavy} ${dest_file} + for op_name in ${name} ; do + add_ops ${op_name} ${isHeavy} ${dest_file} + done done echo "}" >> ${dest_file} file_count=$(cat ${dest_file} | wc -l) -- Gitee