Merge pull request #282 from TESSEorg/hyndavi/feature/readme-update-d…

…evice

update `README` with device example

.github/workflows/cmake.yml

@@ -110,8 +110,10 @@ jobs:
       working-directory: ${{github.workspace}}/build
       shell: bash
       run: |
-        cmake -S $GITHUB_WORKSPACE/doc/dox/dev/devsamp/main -B test_install_devsamp -DCMAKE_PREFIX_PATH=${{github.workspace}}/install || (cat test_install_devsamp/CMakeFiles/CMakeOutput.log && cat test_install_devsamp/CMakeFiles/CMakeError.log)
-        cmake --build test_install_devsamp
+        cmake -S $GITHUB_WORKSPACE/doc/dox/dev/devsamp/main -B test_install_devsamp_main -DCMAKE_PREFIX_PATH=${{github.workspace}}/install || (cat test_install_devsamp_main/CMakeFiles/CMakeOutput.log && cat test_install_devsamp_main/CMakeFiles/CMakeError.log)
+        cmake --build test_install_devsamp_main
+        cmake -S $GITHUB_WORKSPACE/doc/dox/dev/devsamp/fibonacci -B test_install_devsamp_fibonacci -DCMAKE_PREFIX_PATH=${{github.workspace}}/install || (cat test_install_devsamp_fibonacci/CMakeFiles/CMakeOutput.log && cat test_install_devsamp_fibonacci/CMakeFiles/CMakeError.log)
+        cmake --build test_install_devsamp_fibonacci
         cmake -E make_directory test_install_userexamples
         cat > test_install_userexamples/CMakeLists.txt <<EOF
         cmake_minimum_required(VERSION 3.14)

INSTALL.md

@@ -1,6 +1,6 @@
 # synopsis
 
-```.sh
+```sh
 $ git clone https://github.com/TESSEorg/ttg.git
 $ cmake -S ttg -B ttg/build -DCMAKE_INSTALL_PREFIX=/path/to/ttg/install [optional cmake args]
 $ cmake --build ttg/build

doc/dox/dev/devsamp/fibonacci/CMakeLists.txt

@@ -0,0 +1,14 @@
+cmake_minimum_required(VERSION 3.14)
+project(ttg-devsample-fibonacci)
+
+find_package(ttg REQUIRED)
+
+add_ttg_executable(fibonacci fibonacci.cc NOT_EXCLUDE_FROM_ALL)
+# Fib device test
+if (TTG_HAVE_CUDA)
+    add_ttg_executable(fibonacci_cuda
+            fibonacci_device.cc
+            fibonacci_cuda_kernel.h
+            fibonacci_cuda_kernel.cu
+            LINK_LIBRARIES std::coroutine RUNTIMES "parsec" NOT_EXCLUDE_FROM_ALL)
+endif()

doc/dox/dev/devsamp/fibonacci/fibonacci.cc

@@ -0,0 +1,60 @@
+#include <ttg.h>
+#include "ttg/serialization.h"
+
+/// N.B. contains values of F_n and F_{n-1}
+struct Fn {
+  int64_t F[2];  // F[0] = F_n, F[1] = F_{n-1}
+  Fn() { F[0] = 1; F[1] = 0; }
+  template <typename Archive>
+  void serialize(Archive& ar) {
+    ar & F;
+  }
+  template <typename Archive>
+  void serialize(Archive& ar, const unsigned int) {
+    ar & F;
+  }
+};
+auto make_ttg_fib_lt(const int64_t F_n_max = 1000) {
+  ttg::Edge<int64_t, Fn> f2f;
+  ttg::Edge<void, Fn> f2p;
+
+  auto fib = ttg::make_tt(
+      [=](int64_t n, Fn&& f_n) {
+        int64_t next_f_n = f_n.F[0] + f_n.F[1];
+        f_n.F[1] = f_n.F[0];
+        f_n.F[0] = next_f_n;
+        if (next_f_n < F_n_max) {
+          ttg::send<0>(n + 1, f_n);
+        } else {
+          ttg::sendv<1>(f_n);
+        }
+      },
+      ttg::edges(f2f), ttg::edges(f2f, f2p), "fib");
+
+  auto print = ttg::make_tt(
+      [=](Fn&& f_n) {
+        std::cout << "The largest Fibonacci number smaller than " << F_n_max << " is " << f_n.F[1] << std::endl;
+      },
+      ttg::edges(f2p), ttg::edges(), "print");
+  auto ins = std::make_tuple(fib->template in<0>());
+  std::vector<std::unique_ptr<ttg::TTBase>> ops;
+  ops.emplace_back(std::move(fib));
+  ops.emplace_back(std::move(print));
+  return make_ttg(std::move(ops), ins, std::make_tuple(), "Fib_n < N");
+}
+
+int main(int argc, char* argv[]) {
+  ttg::initialize(argc, argv, -1);
+  int64_t N = (argc > 1) ? std::atol(argv[1]) : 1000;
+
+  auto fib = make_ttg_fib_lt(N);
+  ttg::make_graph_executable(fib.get());
+  if (ttg::default_execution_context().rank() == 0)
+    fib->template in<0>()->send(1, Fn{});;
+
+  ttg::execute();
+  ttg::fence();
+
+  ttg::finalize();
+  return 0;
+}

doc/dox/dev/devsamp/fibonacci/fibonacci_cuda_kernel.cu

@@ -0,0 +1,15 @@
+#include "fibonacci_cuda_kernel.h"
+
+#ifdef TTG_HAVE_CUDA
+
+__global__ void cu_next_value(int64_t* fn_and_fnm1) {
+  int64_t fnp1 = fn_and_fnm1[0] + fn_and_fnm1[1];
+  fn_and_fnm1[1] = fn_and_fnm1[0];
+  fn_and_fnm1[0] = fnp1;
+}
+
+void next_value(int64_t* fn_and_fnm1) {
+  cu_next_value<<<1, 1>>>(fn_and_fnm1);
+}
+
+#endif // TTG_HAVE_CUDA

doc/dox/dev/devsamp/fibonacci/fibonacci_cuda_kernel.h

@@ -0,0 +1,4 @@
+#include "ttg/config.h"
+#include <cinttypes>
+
+void next_value(int64_t* fn_and_fnm1);

doc/dox/dev/devsamp/fibonacci/fibonacci_device.cc

@@ -0,0 +1,88 @@
+#include <ttg.h>
+
+#if defined(TTG_HAVE_CUDA)
+#define ES ttg::ExecutionSpace::CUDA
+#include "cuda_runtime.h"
+#include "fibonacci_cuda_kernel.h"
+#else
+#error " CUDA  is required to build this test!"
+#endif
+
+#include "ttg/serialization.h"
+
+const int64_t F_n_max = 1000;
+/// N.B. contains values of F_n and F_{n-1}
+struct Fn : public ttg::TTValue<Fn> {
+  std::unique_ptr<int64_t[]> F;  // F[0] = F_n, F[1] = F_{n-1}
+  ttg::Buffer<int64_t> b;
+
+  Fn() : F(std::make_unique<int64_t[]>(2)), b(F.get(), 2) { F[0] = 1; F[1] = 0; }
+
+  Fn(const Fn&) = delete;
+  Fn(Fn&& other) = default;
+  Fn& operator=(const Fn& other) = delete;
+  Fn& operator=(Fn&& other) = default;
+
+  template <typename Archive>
+  void serialize(Archive& ar) {
+    ttg::ttg_abort();
+  }
+  template <typename Archive>
+  void serialize(Archive& ar, const unsigned int) {
+    ttg::ttg_abort();
+  }
+};
+
+auto make_ttg_fib_lt(const int64_t F_n_max = 1000) {
+  ttg::Edge<int64_t, Fn> f2f;
+  ttg::Edge<void, Fn> f2p;
+
+  auto fib = ttg::make_tt<ES>(
+      [=](int64_t n, Fn&& f_n) -> ttg::device::Task {
+        assert(n > 0);
+        ttg::trace("in fib: n=", n, " F_n=", f_n.F[0]);
+
+        co_await ttg::device::select(f_n.b);
+
+        next_value(f_n.b.current_device_ptr());
+
+        // wait for the task to complete and the values to be brought back to the host
+        co_await ttg::device::wait(f_n.b);
+
+        if (f_n.F[0] < F_n_max) {
+          co_await ttg::device::forward(ttg::device::send<0>(n + 1, std::move(f_n)));
+        } else {
+          co_await ttg::device::forward(ttg::device::sendv<1>(std::move(f_n)));
+        }
+      },
+      ttg::edges(f2f), ttg::edges(f2f, f2p), "fib");
+  auto print = ttg::make_tt(
+      [=](Fn&& f_n) {
+        std::cout << "The largest Fibonacci number smaller than " << F_n_max << " is " << f_n.F[1] << std::endl;
+      },
+      ttg::edges(f2p), ttg::edges(), "print");
+
+  auto ins = std::make_tuple(fib->template in<0>());
+  std::vector<std::unique_ptr<::ttg::TTBase>> ops;
+  ops.emplace_back(std::move(fib));
+  ops.emplace_back(std::move(print));
+  return make_ttg(std::move(ops), ins, std::make_tuple(), "Fib_n < N");
+}
+
+int main(int argc, char* argv[]) {
+  ttg::initialize(argc, argv, -1);
+  ttg::trace_on();
+  int64_t N = 1000;
+  if (argc > 1) N = std::atol(argv[1]);
+  auto fib = make_ttg_fib_lt(N);  // computes largest F_n < N
+
+  ttg::make_graph_executable(fib.get());
+  if (ttg::default_execution_context().rank() == 0)
+    fib->template in<0>()->send(1, Fn{});;
+
+  ttg::execute(ttg::ttg_default_execution_context());
+  ttg::fence(ttg::ttg_default_execution_context());
+
+  ttg::finalize();
+  return 0;
+}

doc/dox/dev/devsamp/main/CMakeLists.txt

@@ -1,5 +1,5 @@
 cmake_minimum_required(VERSION 3.14)
-project(test)
+project(ttg-devsample-main)
 
 find_package(ttg REQUIRED)
 

tests/unit/CMakeLists.txt

@@ -27,11 +27,18 @@ add_ttg_executable(core-unittests-ttg "${ut_src}" LINK_LIBRARIES "${ut_libs}" CO
 add_ttg_executable(serialization serialization.cc unit_main.cpp
                    LINK_LIBRARIES Catch2::Catch2 ttg-serialization  $<TARGET_NAME_IF_EXISTS:BTAS::BTAS>
                    COMPILE_DEFINITIONS $<$<TARGET_EXISTS:BTAS::BTAS>:TTG_HAS_BTAS=1>)
-#target_link_libraries(serialization "Catch2::Catch2;ttg-serialization")
-#if (TARGET BTAS::BTAS)
-#    target_link_libraries(serialization BTAS::BTAS)
-#    target_compile_definitions(serialization PRIVATE TTG_HAS_BTAS=1)
-#endif (TARGET BTAS::BTAS)
+
+# Boost serialization test: checks low-level codegen
+add_ttg_executable(serialization_boost serialization_boost.cc
+        LINK_LIBRARIES ttg-serialization-boost RUNTIMES "parsec")
+
+# Fib device test
+if (TTG_HAVE_CUDA)
+       add_ttg_executable(fibonacci_device fibonacci_device.cc
+            fibonacci_cuda_kernel.h
+            fibonacci_cuda_kernel.cu
+        LINK_LIBRARIES std::coroutine RUNTIMES "parsec")
+endif()
 
 # TODO: convert into unit test
 #if (TARGET MADworld)

tests/unit/fibonacci_cuda_kernel.cu

@@ -0,0 +1,15 @@
+#include "fibonacci_cuda_kernel.h"
+
+#ifdef TTG_HAVE_CUDA
+
+__global__ void cu_next_value(int64_t* fn_and_fnm1) {
+  int64_t fnp1 = fn_and_fnm1[0] + fn_and_fnm1[1];
+  fn_and_fnm1[1] = fn_and_fnm1[0];
+  fn_and_fnm1[0] = fnp1;
+}
+
+void next_value(int64_t* fn_and_fnm1) {
+  cu_next_value<<<1, 1>>>(fn_and_fnm1);
+}
+
+#endif // TTG_HAVE_CUDA

tests/unit/fibonacci_cuda_kernel.h

@@ -0,0 +1,4 @@
+#include "ttg/config.h"
+#include <cinttypes>
+
+void next_value(int64_t* fn_and_fnm1);

tests/unit/fibonacci_device.cc

@@ -0,0 +1,88 @@
+#include <ttg.h>
+
+#if defined(TTG_HAVE_CUDA)
+#define ES ttg::ExecutionSpace::CUDA
+#include "cuda_runtime.h"
+#include "fibonacci_cuda_kernel.h"
+#else
+#error " CUDA  is required to build this test!"
+#endif
+
+#include "ttg/serialization.h"
+
+const int64_t F_n_max = 1000;
+/// N.B. contains values of F_n and F_{n-1}
+struct Fn : public ttg::TTValue<Fn> {
+  std::unique_ptr<int64_t[]> F;  // F[0] = F_n, F[1] = F_{n-1}
+  ttg::Buffer<int64_t> b;
+
+  Fn() : F(std::make_unique<int64_t[]>(2)), b(F.get(), 2) { F[0] = 1; F[1] = 0; }
+
+  Fn(const Fn&) = delete;
+  Fn(Fn&& other) = default;
+  Fn& operator=(const Fn& other) = delete;
+  Fn& operator=(Fn&& other) = default;
+
+  template <typename Archive>
+  void serialize(Archive& ar) {
+    ttg::ttg_abort();
+  }
+  template <typename Archive>
+  void serialize(Archive& ar, const unsigned int) {
+    ttg::ttg_abort();
+  }
+};
+
+auto make_ttg_fib_lt(const int64_t F_n_max = 1000) {
+  ttg::Edge<int64_t, Fn> f2f;
+  ttg::Edge<void, Fn> f2p;
+
+  auto fib = ttg::make_tt<ES>(
+      [=](int64_t n, Fn&& f_n) -> ttg::device::Task {
+        assert(n > 0);
+        ttg::trace("in fib: n=", n, " F_n=", f_n.F[0]);
+
+        co_await ttg::device::select(f_n.b);
+
+        next_value(f_n.b.current_device_ptr());
+
+        // wait for the task to complete and the values to be brought back to the host
+        co_await ttg::device::wait(f_n.b);
+
+        if (f_n.F[0] < F_n_max) {
+          co_await ttg::device::forward(ttg::device::send<0>(n + 1, std::move(f_n)));
+        } else {
+          co_await ttg::device::forward(ttg::device::sendv<1>(std::move(f_n)));
+        }
+      },
+      ttg::edges(f2f), ttg::edges(f2f, f2p), "fib");
+  auto print = ttg::make_tt(
+      [=](Fn&& f_n) {
+        std::cout << "The largest Fibonacci number smaller than " << F_n_max << " is " << f_n.F[1] << std::endl;
+      },
+      ttg::edges(f2p), ttg::edges(), "print");
+
+  auto ins = std::make_tuple(fib->template in<0>());
+  std::vector<std::unique_ptr<::ttg::TTBase>> ops;
+  ops.emplace_back(std::move(fib));
+  ops.emplace_back(std::move(print));
+  return make_ttg(std::move(ops), ins, std::make_tuple(), "Fib_n < N");
+}
+
+int main(int argc, char* argv[]) {
+  ttg::initialize(argc, argv, -1);
+  ttg::trace_on();
+  int64_t N = 1000;
+  if (argc > 1) N = std::atol(argv[1]);
+  auto fib = make_ttg_fib_lt(N);  // computes largest F_n < N
+
+  ttg::make_graph_executable(fib.get());
+  if (ttg::default_execution_context().rank() == 0)
+    fib->template in<0>()->send(1, Fn{});;
+
+  ttg::execute(ttg::ttg_default_execution_context());
+  ttg::fence(ttg::ttg_default_execution_context());
+
+  ttg::finalize();
+  return 0;
+}