From cfb010b6c7a03eeba59b09af56183b8cc3b083ed Mon Sep 17 00:00:00 2001 From: Ronan Keryell Date: Tue, 27 Feb 2024 19:16:20 -0800 Subject: [PATCH] Add attributes for g++ and clang++ to export the emulator functions Be sure the weak functions are exported to be linked with MLIR-AIE output. --- src/CMakeLists.txt | 1 + src/mlir-aie-emulator/CMakeLists.txt | 2 ++ src/mlir-aie-emulator/mlir-aie-emulator.cpp | 39 ++++++++++++++++----- 3 files changed, 34 insertions(+), 8 deletions(-) create mode 100644 src/CMakeLists.txt create mode 100644 src/mlir-aie-emulator/CMakeLists.txt diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt new file mode 100644 index 00000000..8b3944c7 --- /dev/null +++ b/src/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(mlir-aie-emulator) diff --git a/src/mlir-aie-emulator/CMakeLists.txt b/src/mlir-aie-emulator/CMakeLists.txt new file mode 100644 index 00000000..c5c9e1ac --- /dev/null +++ b/src/mlir-aie-emulator/CMakeLists.txt @@ -0,0 +1,2 @@ +add_executable(mlir-aie-emulator mlir-aie-emulator.cpp) +add_sycl_to_target(mlir-aie-emulator) diff --git a/src/mlir-aie-emulator/mlir-aie-emulator.cpp b/src/mlir-aie-emulator/mlir-aie-emulator.cpp index 53a9fbfa..270b1b5a 100644 --- a/src/mlir-aie-emulator/mlir-aie-emulator.cpp +++ b/src/mlir-aie-emulator/mlir-aie-emulator.cpp @@ -10,40 +10,63 @@ #include #include +// The compiler-dependent attribute to avoid the compiler to optimized too much +#ifdef __clang__ +// Compiled with clang++ +#define EMULATOR_CONFUSING_ATTRIBUTE noinline +// This macro is also defined by clang++, so test clang++ first +#elif defined(__GNUC__) +// Compiled with g++ +#define EMULATOR_CONFUSING_ATTRIBUTE noipa +#else +#error Compiler not handled +#endif + namespace aie = sycl::vendor::xilinx::acap::aie; /// There is a G++ bug where weak templated symbols do not get emitted when /// optimizing. Use this function to pass a symbol too so the compiler looses /// track of the use and really emits the symbol. Inspired by /// https://gcc.gnu.org/bugzilla/show_bug.cgi?id=61882#c2 -__attribute__((weak)) __attribute__((noinline)) void +__attribute__((weak)) __attribute__((EMULATOR_CONFUSING_ATTRIBUTE)) void aie_tile_confuse_the_compiler(auto value) {} namespace { /// To avoid passing a tile handle though all the MLIR generated code, just put /// the handle in a global variable which is used by some trampoline functions -template This aie_tile_handle; +template This aie_tile_handle; /// A dispatch function to write to the output port -template +template void (*aie_tile_put_ms_trampoline)(std::int32_t port, std::int32_t value); } // namespace /// A dispatch function to write to an output port -template +/// +/// This is mangled by clang++ for Col=1 and Row=3 as +/// _Z25aie_tile_put_ms_intrinsicITnDaLi1ETnDaLi3EEvii +/// +/// This is mangled by g++ for Col=1 and Row=3 as +/// _Z25aie_tile_put_ms_intrinsicILi1ELi3EEvii +template __attribute__((weak)) __attribute__((noinline)) void aie_tile_put_ms_intrinsic(std::int32_t port, std::int32_t value) { - aie_tile_put_ms_trampoline(port, value); + aie_tile_put_ms_trampoline(port, value); } /// Use some weak symbol function so the linker can replace this empty function /// by any MLIR-generated tile program function /// -/// This is mangled by G++ for as _Z16aie_tile_programILi1ELi1EEvv -template -__attribute__((weak)) __attribute__((noinline)) void aie_tile_program() {} +/// This is mangled by clang++ for Col=1 and Row=3 as +/// _Z16aie_tile_programITnDaLi1ETnDaLi3EEvv +/// +/// This is mangled by g++ for Col=1 and Row=3 as +/// _Z16aie_tile_programILi1ELi3EEvv +template +__attribute__((weak)) __attribute__((EMULATOR_CONFUSING_ATTRIBUTE)) void +aie_tile_program() {} namespace {