From 2573c9e25096da87cc87bd6305d8c869409db09c Mon Sep 17 00:00:00 2001 From: Ronan Keryell Date: Mon, 26 Feb 2024 18:42:52 -0800 Subject: [PATCH] Add runtime prototype to emulate MLIR-AIE programs --- include/triSYCL/detail/global_config.hpp | 3 +- src/mlir-aie-emulator/mlir-aie-emulator.cpp | 109 ++++++++++---------- 2 files changed, 58 insertions(+), 54 deletions(-) diff --git a/include/triSYCL/detail/global_config.hpp b/include/triSYCL/detail/global_config.hpp index 1d096412c..39038b535 100644 --- a/include/triSYCL/detail/global_config.hpp +++ b/include/triSYCL/detail/global_config.hpp @@ -110,7 +110,8 @@ namespace trisycl::detail { /// @} End the defaults Doxygen group -// Compiler specific weak linking (until changing to C++17 inline variables/functions) +// Compiler specific weak linking (mostly until changing to C++17 inline +// variables/functions) #ifndef TRISYCL_WEAK_ATTRIB_PREFIX #ifdef _MSC_VER #define TRISYCL_WEAK_ATTRIB_PREFIX __declspec(selectany) diff --git a/src/mlir-aie-emulator/mlir-aie-emulator.cpp b/src/mlir-aie-emulator/mlir-aie-emulator.cpp index f75a1e705..53a9fbfaf 100644 --- a/src/mlir-aie-emulator/mlir-aie-emulator.cpp +++ b/src/mlir-aie-emulator/mlir-aie-emulator.cpp @@ -1,9 +1,8 @@ -/* Simple executor for MLIR-generated AIR on AIE using ACAP++ +/// Simple executor for MLIR-AIE code based on ACAP++ +/// +/// RUN: %{execute}%s - RUN: %{execute}%s -*/ - -// Put the tile code on fiber too to boost the performances +/// Put the tile code on fiber too to boost the performances #define TRISYCL_XILINX_AIE_TILE_CODE_ON_FIBER 1 #include @@ -11,68 +10,72 @@ #include #include -using namespace sycl::vendor::xilinx; -using namespace sycl::vendor::xilinx::acap::aie; - -template void (*aie_tile_out_intrinsic)(std::int32_t port, std::int32_t value); - -/** Use some weak symbol function so the linker can replace it by a - more interesting one -*/ -template -TRISYCL_WEAK_ATTRIB_PREFIX void TRISYCL_WEAK_ATTRIB_SUFFIX -air_tile(void* self, void (*out0_write)(void* self, unsigned int value)) { - /* By linking with some AIR code such as _Z8air_tileILi6ELi4EEvPvPFvS0_jE - (g++) or _Z8air_tileITnDaLi6ETnDaLi4EEvPvPFvS0_jE (clang) for air_tile for - X=6 and Y=4, this weak symbol will be replaced by the provided function - instead. - */ - // To write some value to out 0 port - out0_write(self, 42); -} +namespace aie = sycl::vendor::xilinx::acap::aie; -/** Use some weak symbol function so the linker can replace it by a - more interesting one -*/ -template -TRISYCL_WEAK_ATTRIB_PREFIX void TRISYCL_WEAK_ATTRIB_SUFFIX -aie_tile() { -} +/// 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 +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; -/* -auto type_erase(auto function) { +/// A dispatch function to write to the output port +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 +__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); } -*/ -/** An executor kernel just calling some AIR code and exporting some - useful functions */ -template -struct executor : acap::aie::tile { - using t = acap::aie::tile; + +/// 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() {} + +namespace { + +/// An executor kernel just calling some AIR code and exporting some useful +/// functions and variables to be used +template struct executor : aie::tile { + using t = aie::tile; void run() { - // Export like this any useful ACAP++ function needed by your code - auto out0_write = [](void* self, unsigned int value) { - reinterpret_cast(self)->out(0) << value; - }; -// std::array dispatch_table { type_erase(+) }; - // Call the potential AIR tile code - air_tile(this, out0_write); - aie_tile_out_intrinsic = [this](std::int32_t port, - std::int32_t value) { - out(port) << value; - }; - aie_tile(); + // Save "this" into global variable to avoid capturing it in the lambda or + // to have to pass it through all the user code + aie_tile_handle = this; + aie_tile_put_ms_trampoline = + +[](std::int32_t port, std::int32_t value) { + aie_tile_handle->out(port) << value; + }; + // Run the tile program + aie_tile_program(); + /// Just reference the intrinsic functions to be sure the functions get + /// instantiated + aie_tile_confuse_the_compiler(aie_tile_put_ms_intrinsic); } }; +} // namespace int main() try { // acap::aie::device {}.run(); - acap::aie::device> {}.run(); + aie::device> {}.run(); } catch (sycl::exception& e) { // Display the string message of any SYCL exception std::cerr << e.what() << std::endl; - // Rethrow to make clear something bad happened + // Rethrow to make it clear that something bad happened throw; }