Add CINN Compile Option (#36292)

Add CINN compile option in CMake.

Now you can use CINN in Paddle by `-DWITH_CINN=ON` when `cmake`

To test it, you can run `make cinn_lib_test -j` and `ctest -R cinn_lib_test`. 

Note:
1. You should set
```
export runtime_include_dir=${CINN_SOURCE_DIR}/cinn/runtime/cuda 
```
When run test, the `${CINN_SOURCE_DIR}` should be set based on your CINN directory.

2. CINN is under developing now, you may have to change `CINN_GIT_TAG` to the git commit you need.

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

cmake/cinn.cmake

+# 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)
+

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)

paddle/fluid/framework/ir/cinn_lib_test.cc

+/* 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 <glog/logging.h>
+#include <gtest/gtest.h>
+
+#include <algorithm>
+#include <memory>
+#include <random>
+#include <vector>
+
+#ifdef PADDLE_WITH_CUDA
+#include <cuda_runtime.h>
+#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<float>(target);
+  std::random_device seed;
+  std::default_random_engine engine(seed());
+  std::uniform_real_distribution<float> dist(0.f, 1.f);
+  size_t num_ele = tensor->shape().numel();
+  std::vector<float> 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<hlir::framework::Graph>(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<hlir::framework::Tensor>("A");
+  scope->Var<hlir::framework::Tensor>("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<hlir::framework::Graph>(program, target);
+  auto scope = BuildScope(target, graph);
+  hlir::framework::GraphCompiler gc(target, scope, graph);
+  auto runtime_program = gc.Build();
+
+  scope->Var<hlir::framework::Tensor>(std::string(a.id()));
+  scope->Var<hlir::framework::Tensor>(std::string(w.id()));
+  scope->Var<hlir::framework::Tensor>(std::string(b.id()));
+  scope->Var<hlir::framework::Tensor>(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