diff --git a/CMakeLists.txt b/CMakeLists.txt index 98772e967815314ec0603aab5363202ec5e21713..d4a0eb067b4f17762142fc41ea2f2b0dc4207278 100755 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -214,6 +214,7 @@ option(PY_VERSION "Compile PaddlePaddle with python3 support" ${PY_VER option(WITH_DGC "Use DGC(Deep Gradient Compression) or not" ${WITH_DISTRIBUTE}) option(SANITIZER_TYPE "Choose the type of sanitizer, options are: Address, Leak, Memory, Thread, Undefined" OFF) option(WITH_LITE "Compile Paddle Fluid with Lite Engine" OFF) +option(WITH_CINN "Compile PaddlePaddle with CINN" OFF) option(WITH_NCCL "Compile PaddlePaddle with NCCL support" ON) option(WITH_RCCL "Compile PaddlePaddle with RCCL support" ON) option(WITH_XPU_BKCL "Compile PaddlePaddle with BAIDU KUNLUN XPU BKCL" OFF) @@ -299,6 +300,10 @@ if(WITH_GPU) endif() endif() +if(WITH_CINN) + include(cinn) +endif() + if(WITH_ROCM) include(hip) include(miopen) # set miopen libraries, must before configure diff --git a/cmake/cinn.cmake b/cmake/cinn.cmake new file mode 100644 index 0000000000000000000000000000000000000000..dd5f809e9581a211f585eb663b782caa8e857e81 --- /dev/null +++ b/cmake/cinn.cmake @@ -0,0 +1,112 @@ +# Copyright (c) 2021 PaddlePaddle Authors. 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. + +if (NOT WITH_CINN) + return() +endif() + +# TODO(zhhsplendid): CINN has lots of warnings during early development. +# They will be treated as errors under paddle. We set no-error now and we will +# clean the code in the future. +add_definitions(-w) + +###################################### +# Build CINN from Git External Project +###################################### +include(ExternalProject) +set(CINN_SOURCE_DIR ${THIRD_PARTY_PATH}/CINN) +# TODO(zhhsplendid): Modify git tag after we have release tag +set(CINN_GIT_TAG 3f004bfa3ed273ecf1de8e7b946433038c79b84f) +set(CINN_OPTIONAL_ARGS -DWITH_CUDA=${WITH_GPU} -DWITH_CUDNN=${WITH_GPU} -DPUBLISH_LIBS=ON) +set(CINN_BUILD_COMMAND $(MAKE) cinncore -j && $(MAKE) cinnapi -j) +ExternalProject_Add( + external_cinn + ${EXTERNAL_PROJECT_LOG_ARGS} + GIT_REPOSITORY "${GIT_URL}/PaddlePaddle/CINN.git" + GIT_TAG ${CINN_GIT_TAG} + PREFIX ${CINN_SOURCE_DIR} + UPDATE_COMMAND "" + BUILD_COMMAND ${CINN_BUILD_COMMAND} + INSTALL_COMMAND "" + CMAKE_ARGS ${CINN_OPTIONAL_ARGS}) + + + +ExternalProject_Get_property(external_cinn BINARY_DIR) +ExternalProject_Get_property(external_cinn SOURCE_DIR) +set(CINN_BINARY_DIR ${BINARY_DIR}) +set(CINN_SOURCE_DIR ${SOURCE_DIR}) + +message(STATUS "CINN BINARY_DIR: ${CINN_BINARY_DIR}") +message(STATUS "CINN SOURCE_DIR: ${CINN_SOURCE_DIR}") + + +######################### +# Add CINN's dependencies +######################### + +# Add absl +set(ABSL_LIB_NAMES + hash + wyhash + city + strings + throw_delegate + bad_any_cast_impl + bad_optional_access + bad_variant_access + raw_hash_set + ) +set(ABSL_LIB_DIR "${CINN_BINARY_DIR}/dist/third_party/absl/lib") +set(ABSL_INCLUDE_DIR "${CINN_BINARY_DIR}/dist/third_party/absl/include") +add_library(absl STATIC IMPORTED GLOBAL) +set_target_properties(absl PROPERTIES IMPORTED_LOCATION ${ABSL_LIB_DIR}/libabsl_base.a) +foreach(lib_name ${ABSL_LIB_NAMES}) + target_link_libraries(absl INTERFACE ${ABSL_LIB_DIR}/libabsl_${lib_name}.a) +endforeach() +include_directories(${ABSL_INCLUDE_DIR}) + +# Add isl +set(ISL_LIB_DIR "${CINN_BINARY_DIR}/dist/third_party/isl/lib") +set(ISL_INCLUDE_DIR "${CINN_BINARY_DIR}/dist/third_party/isl/include") +add_library(isl STATIC IMPORTED GLOBAL) +set_target_properties(isl PROPERTIES IMPORTED_LOCATION ${ISL_LIB_DIR}/libisl.a) +include_directories(${ISL_INCLUDE_DIR}) + +# Add LLVM +set(LLVM_LIB_NAMES + ExecutionEngine + ) +set(LLVM_LIB_DIR "${CINN_BINARY_DIR}/dist/third_party/llvm/lib") +set(LLVM_INCLUDE_DIR "${CINN_BINARY_DIR}/dist/third_party/llvm/include") +add_library(llvm STATIC IMPORTED GLOBAL) +set_target_properties(llvm PROPERTIES IMPORTED_LOCATION ${LLVM_LIB_DIR}/libLLVMCore.a) +foreach(lib_name ${LLVM_LIB_NAMES}) + target_link_libraries(llvm INTERFACE ${LLVM_LIB_DIR}/libLLVM${lib_name}.a) +endforeach() +include_directories(${LLVM_INCLUDE_DIR}) + +###################################################### +# Put external_cinn and dependencies together as a lib +###################################################### + +set(CINN_LIB_NAME "libcinnapi.so") +set(CINN_LIB_LOCATION "${CINN_BINARY_DIR}/dist/cinn/lib") +set(CINN_INCLUDE_DIR "${CINN_BINARY_DIR}/dist/cinn/include") + +add_library(cinn SHARED IMPORTED GLOBAL) +set_target_properties(cinn PROPERTIES IMPORTED_LOCATION "${CINN_LIB_LOCATION}/${CINN_LIB_NAME}") +include_directories(${CINN_INCLUDE_DIR}) +add_dependencies(cinn external_cinn absl isl llvm glog gflag) + diff --git a/paddle/fluid/framework/ir/CMakeLists.txt b/paddle/fluid/framework/ir/CMakeLists.txt index 904450b5b251ee26c80f051d8fa945302329b3c2..7b80d331ff7077c526320573ebce33c48036a2c6 100644 --- a/paddle/fluid/framework/ir/CMakeLists.txt +++ b/paddle/fluid/framework/ir/CMakeLists.txt @@ -143,6 +143,9 @@ cc_test(pass_test SRCS pass_test.cc DEPS graph pass graph_helper) cc_test(graph_test SRCS graph_test.cc DEPS graph graph_helper op_registry) cc_test(graph_helper_test SRCS graph_helper_test.cc DEPS graph graph_helper op_registry) cc_test(graph_to_program_pass_test SRCS graph_to_program_pass_test.cc DEPS graph_to_program_pass) +if (WITH_CINN) + cc_test(cinn_lib_test SRCS cinn_lib_test.cc DEPS cinn) +endif() cc_test(cost_model_test SRCS cost_model_test.cc DEPS cost_model op_registry) cc_test(test_graph_pattern_detector SRCS graph_pattern_detector_tester.cc DEPS graph_pattern_detector) cc_test(test_op_compat_sensible_pass SRCS op_compat_sensible_pass_tester.cc DEPS op_compat_sensible_pass) diff --git a/paddle/fluid/framework/ir/cinn_lib_test.cc b/paddle/fluid/framework/ir/cinn_lib_test.cc new file mode 100644 index 0000000000000000000000000000000000000000..cdee45a06c71afb86d3b7577190f3bd8a7eca340 --- /dev/null +++ b/paddle/fluid/framework/ir/cinn_lib_test.cc @@ -0,0 +1,151 @@ +/* Copyright (c) 2021 PaddlePaddle Authors. 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 +#include +#include +#include + +#ifdef PADDLE_WITH_CUDA +#include +#endif + +#include "cinn/common/target.h" +#include "cinn/frontend/net_builder.h" +#include "cinn/frontend/syntax.h" +#include "cinn/hlir/framework/graph.h" +#include "cinn/hlir/framework/graph_compiler.h" +#include "cinn/hlir/framework/pass.h" +#include "cinn/hlir/framework/tensor.h" +#include "cinn/hlir/op/use_ops.h" +#include "cinn/hlir/pass/use_pass.h" + +namespace cinn { +namespace frontend { + +Program CreateAddProgram() { + constexpr int M = 32; + constexpr int N = 24; + + NetBuilder builder("net_builder"); + auto a = builder.CreateInput(Float(32), {M, N}); + auto b = builder.CreateInput(Float(32), {M, N}); + auto c = builder.add(a, b); + auto d = builder.add(a, c); + auto program = builder.Build(); + + return program; +} + +void SetRandData(hlir::framework::Tensor tensor, Target target) { + auto* data = tensor->mutable_data(target); + std::random_device seed; + std::default_random_engine engine(seed()); + std::uniform_real_distribution dist(0.f, 1.f); + size_t num_ele = tensor->shape().numel(); + std::vector random_data(num_ele); + for (size_t i = 0; i < num_ele; i++) { + random_data[i] = dist(engine); // All random data + } + +#ifdef PADDLE_WITH_CUDA + cudaMemcpy(data, random_data.data(), num_ele * sizeof(float), + cudaMemcpyHostToDevice); +#else + std::copy(random_data.begin(), random_data.end(), data); +#endif +} + +TEST(net_build, basic) { + auto program = CreateAddProgram(); + // output program + for (size_t i = 0; i < program.size(); i++) { + LOG(INFO) << "instruction: " << program[i]; + } +} + +TEST(net_build, program_execute_multi_elementwise_add) { + auto program = CreateAddProgram(); +#ifdef PADDLE_WITH_CUDA + Target target = common::DefaultNVGPUTarget(); +#else + Target target = common::DefaultHostTarget(); +#endif + + auto graph = std::make_shared(program, target); + std::cout << "graph:\n" << graph->Visualize() << std::endl; + + auto scope = BuildScope(target, graph); + hlir::framework::GraphCompiler gc(target, scope, graph); + auto runtime_program = gc.Build(); + + scope->Var("A"); + scope->Var("B"); + + auto A = scope->GetTensor("A"); + auto B = scope->GetTensor("B"); + SetRandData(A, target); + SetRandData(B, target); + + runtime_program->Execute(); +} + +TEST(net_build, program_execute_fc) { + constexpr int B = 10; // batch size + constexpr int M = 32; + constexpr int K = 18; + constexpr int N = 24; + + NetBuilder builder("net_builder"); + auto a = builder.CreateInput(Float(32), {B, M, K}, "A"); + auto w = builder.CreateInput(Float(32), {N, K}, "W"); // weight + auto b = builder.CreateInput(Float(32), {N}, "B"); // bias + + auto mul_out = builder.mul(a, w, 2, 1); + auto add_out = builder.add(mul_out, b); + auto program = builder.Build(); + +#ifdef PADDLE_WITH_CUDA + Target target = common::DefaultNVGPUTarget(); +#else + Target target = common::DefaultHostTarget(); +#endif + + auto graph = std::make_shared(program, target); + auto scope = BuildScope(target, graph); + hlir::framework::GraphCompiler gc(target, scope, graph); + auto runtime_program = gc.Build(); + + scope->Var(std::string(a.id())); + scope->Var(std::string(w.id())); + scope->Var(std::string(b.id())); + scope->Var(std::string(mul_out->id)); + + auto a_ten = scope->GetTensor(std::string(a.id())); + auto w_ten = scope->GetTensor(std::string(w.id())); + auto b_ten = scope->GetTensor(std::string(b.id())); + auto fake_out_ten = scope->GetTensor(std::string(mul_out->id)); + auto add_out_ten = scope->GetTensor(std::string(add_out->id)); + SetRandData(a_ten, target); + SetRandData(w_ten, target); + SetRandData(b_ten, target); + + runtime_program->Execute(); +} + +} // namespace frontend +} // namespace cinn