Implement Kokkos::printf (kokkos#6083)

* Implement Kokkos::printf * Use Kokkos::printf * Update core/unit_test/TestPrintf.hpp Co-authored-by: Damien L-G <dalg24+github@gmail.com> --------- Co-authored-by: Damien L-G <dalg24+github@gmail.com>
nliber · Jul 8, 2023 · 597bc36 · 597bc36
1 parent 53a5aef
commit 597bc36
Show file tree

Hide file tree

Showing 17 changed files with 208 additions and 144 deletions.
diff --git a/containers/src/Kokkos_DynRankView.hpp b/containers/src/Kokkos_DynRankView.hpp
@@ -221,7 +221,7 @@ KOKKOS_INLINE_FUNCTION bool dyn_rank_view_verify_operator_bounds(
     return (size_t(i) < map.extent(R)) &&
            dyn_rank_view_verify_operator_bounds<R + 1>(rank, map, args...);
   } else if (i != 0) {
-    KOKKOS_IMPL_DO_NOT_USE_PRINTF(
+    Kokkos::printf(
         "DynRankView Debug Bounds Checking Error: at rank %u\n  Extra "
         "arguments beyond the rank must be zero \n",
         R);

diff --git a/containers/src/Kokkos_UnorderedMap.hpp b/containers/src/Kokkos_UnorderedMap.hpp
@@ -568,7 +568,7 @@ class UnorderedMap {
           // Previously claimed an unused entry that was not inserted.
           // Release this unused entry immediately.
           if (!m_available_indexes.reset(new_index)) {
-            KOKKOS_IMPL_DO_NOT_USE_PRINTF("Unable to free existing\n");
+            Kokkos::printf("Unable to free existing\n");
           }
         }
 

diff --git a/containers/src/impl/Kokkos_UnorderedMap_impl.hpp b/containers/src/impl/Kokkos_UnorderedMap_impl.hpp
@@ -226,8 +226,8 @@ struct UnorderedMapPrint {
     uint32_t list = m_map.m_hash_lists(i);
     for (size_type curr = list, ii = 0; curr != invalid_index;
          curr = m_map.m_next_index[curr], ++ii) {
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d[%d]: %d->%d\n", list, ii,
-                                    m_map.key_at(curr), m_map.value_at(curr));
+      Kokkos::printf("%d[%d]: %d->%d\n", list, ii, m_map.key_at(curr),
+                     m_map.value_at(curr));
     }
   }
 };

diff --git a/core/src/Kokkos_Core_fwd.hpp b/core/src/Kokkos_Core_fwd.hpp
@@ -27,6 +27,7 @@
 
 #include <Kokkos_Macros.hpp>
 #include <impl/Kokkos_Error.hpp>
+#include <impl/Kokkos_Printf.hpp>
 #include <impl/Kokkos_Utilities.hpp>
 
 #ifdef KOKKOS_ENABLE_DEPRECATED_CODE_3

diff --git a/core/src/Kokkos_Macros.hpp b/core/src/Kokkos_Macros.hpp
@@ -342,7 +342,7 @@
 // Temporary solution for SYCL not supporting printf in kernels.
 // Might disappear at any point once we have found another solution.
 #if !defined(KOKKOS_IMPL_DO_NOT_USE_PRINTF)
-#define KOKKOS_IMPL_DO_NOT_USE_PRINTF(...) printf(__VA_ARGS__)
+#define KOKKOS_IMPL_DO_NOT_USE_PRINTF(...) ::printf(__VA_ARGS__)
 #endif
 
 //----------------------------------------------------------------------------

diff --git a/core/src/Kokkos_ScratchSpace.hpp b/core/src/Kokkos_ScratchSpace.hpp
@@ -126,7 +126,7 @@ class ScratchMemorySpace {
       // mfh 23 Jun 2015: printf call consumes 25 registers
       // in a CUDA build, so only print in debug mode.  The
       // function still returns nullptr if not enough memory.
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
+      Kokkos::printf(
           "ScratchMemorySpace<...>::get_shmem: Failed to allocate "
           "%ld byte(s); remaining capacity is %ld byte(s)\n",
           long(size), long(capacity));

diff --git a/core/src/SYCL/Kokkos_SYCL_Abort.hpp b/core/src/SYCL/Kokkos_SYCL_Abort.hpp
@@ -17,7 +17,7 @@
 #ifndef KOKKOS_SYCL_ABORT_HPP
 #define KOKKOS_SYCL_ABORT_HPP
 
-#include <Kokkos_Macros.hpp>
+#include <impl/Kokkos_Printf.hpp>
 #if defined(KOKKOS_ENABLE_SYCL)
 // FIXME_SYCL
 #if __has_include(<sycl/sycl.hpp>)
@@ -31,7 +31,7 @@ namespace Impl {
 
 inline void sycl_abort(char const* msg) {
 #ifdef NDEBUG
-  KOKKOS_IMPL_DO_NOT_USE_PRINTF("Aborting with message %s.\n", msg);
+  Kokkos::printf("Aborting with message %s.\n", msg);
 #else
   // Choosing "" here causes problems but a single whitespace character works.
   const char* empty = " ";

diff --git a/core/src/impl/Kokkos_Printf.hpp b/core/src/impl/Kokkos_Printf.hpp
@@ -0,0 +1,51 @@
+//@HEADER
+// ************************************************************************
+//
+//                        Kokkos v. 4.0
+//       Copyright (2022) National Technology & Engineering
+//               Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#ifndef KOKKOS_IMPL_PRINTF_HPP
+#define KOKKOS_IMPL_PRINTF_HPP
+
+#include <Kokkos_Macros.hpp>
+
+#ifdef KOKKOS_ENABLE_SYCL
+#include <sycl/sycl.hpp>
+#else
+#include <cstdio>
+#endif
+
+namespace Kokkos {
+
+// In contrast to std::printf, return void to get a consistent behavior across
+// backends. The GPU backends always return 1 and NVHPC only compiles if we
+// don't ask for the return value.
+template <typename... Args>
+KOKKOS_FUNCTION void printf(const char* format, Args... args) {
+#ifdef KOKKOS_ENABLE_SYCL
+  // Some compilers warn if "args" is empty and format is not a string literal
+  if constexpr (sizeof...(Args) == 0)
+    sycl::ext::oneapi::experimental::printf("%s", format);
+  else
+    sycl::ext::oneapi::experimental::printf(format, args...);
+#else
+  if constexpr (sizeof...(Args) == 0)
+    ::printf("%s", format);
+  else
+    ::printf(format, args...);
+#endif
+}
+
+}  // namespace Kokkos
+
+#endif /* #ifndef KOKKOS_IMPL_PRINTF_HPP */
diff --git a/core/unit_test/CMakeLists.txt b/core/unit_test/CMakeLists.txt
@@ -175,6 +175,7 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;OpenMPTarget;OpenACC;HIP;SYCL)
         NumericTraits
         Other
         ParallelScanRangePolicy
+        Printf
         QuadPrecisionMath
         RangePolicy
         RangePolicyConstructors

diff --git a/core/unit_test/TestBitManipulationBuiltins.hpp b/core/unit_test/TestBitManipulationBuiltins.hpp
@@ -77,9 +77,9 @@ struct TestBitManipFunction {
   KOKKOS_FUNCTION void operator()(int i, int& e) const {
     if (Func::eval_builtin(val_[i]) != Func::eval_constexpr(val_[i])) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "value at %x which is %d was expected to be %d\n", (unsigned)val_[i],
-          (int)Func::eval_builtin(val_[i]), (int)Func::eval_constexpr(val_[i]));
+      Kokkos::printf("value at %x which is %d was expected to be %d\n",
+                     (unsigned)val_[i], (int)Func::eval_builtin(val_[i]),
+                     (int)Func::eval_constexpr(val_[i]));
     }
   }
 };
@@ -549,7 +549,7 @@ struct TestBitRotateFunction {
     if (Func::eval_builtin(val_[i].x, val_[i].s) !=
         Func::eval_constexpr(val_[i].x, val_[i].s)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
+      Kokkos::printf(
           "value at %x rotated by %d which is %x was expected to be %x\n",
           (unsigned)val_[i].x, val_[i].s,
           (unsigned)Func::eval_builtin(val_[i].x, val_[i].s),
@@ -726,11 +726,10 @@ struct TestByteswapFunction {
     using Kokkos::Experimental::byteswap_builtin;
     if (byteswap_builtin(value) != expected) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "value at %llx which is %llx was expected to be %llx\n",
-          (unsigned long long)value,
-          (unsigned long long)byteswap_builtin(value),
-          (unsigned long long)expected);
+      Kokkos::printf("value at %llx which is %llx was expected to be %llx\n",
+                     (unsigned long long)value,
+                     (unsigned long long)byteswap_builtin(value),
+                     (unsigned long long)expected);
     }
   }
 };

diff --git a/core/unit_test/TestMathematicalFunctions.hpp b/core/unit_test/TestMathematicalFunctions.hpp
@@ -225,9 +225,8 @@ struct FloatingPointComparison {
 
     bool ar = absolute(fpv) < abs_tol;
     if (!ar) {
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "absolute value exceeds tolerance [|%e| > %e]\n", (double)fpv,
-          abs_tol);
+      Kokkos::printf("absolute value exceeds tolerance [|%e| > %e]\n",
+                     (double)fpv, abs_tol);
     }
 
     return ar;
@@ -248,9 +247,8 @@ struct FloatingPointComparison {
       double rel_diff = abs_diff / min_denom;
       bool ar         = abs_diff == 0 || rel_diff < rel_tol;
       if (!ar) {
-        KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-            "relative difference exceeds tolerance [%e > %e]\n",
-            (double)rel_diff, rel_tol);
+        Kokkos::printf("relative difference exceeds tolerance [%e > %e]\n",
+                       (double)rel_diff, rel_tol);
       }
 
       return ar;
@@ -488,9 +486,9 @@ struct TestMathUnaryFunction : FloatingPointComparison {
     bool ar = compare(Func::eval(val_[i]), res_[i], Func::ulp_factor());
     if (!ar) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "value at %f which is %f was expected to be %f\n", (double)val_[i],
-          (double)Func::eval(val_[i]), (double)res_[i]);
+      Kokkos::printf("value at %f which is %f was expected to be %f\n",
+                     (double)val_[i], (double)Func::eval(val_[i]),
+                     (double)res_[i]);
     }
   }
 };
@@ -533,9 +531,9 @@ struct TestMathBinaryFunction : FloatingPointComparison {
     bool ar = compare(Func::eval(val1_, val2_), res_, Func::ulp_factor());
     if (!ar) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "value at %f, %f which is %f was expected to be %f\n", (double)val1_,
-          (double)val2_, (double)Func::eval(val1_, val2_), (double)res_);
+      Kokkos::printf("value at %f, %f which is %f was expected to be %f\n",
+                     (double)val1_, (double)val2_,
+                     (double)Func::eval(val1_, val2_), (double)res_);
     }
   }
 };
@@ -574,10 +572,9 @@ struct TestMathTernaryFunction : FloatingPointComparison {
         compare(Func::eval(val1_, val2_, val3_), res_, Func::ulp_factor());
     if (!ar) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "value at %f, %f, %f which is %f was expected to be %f\n",
-          (double)val1_, (double)val2_, (double)val3_,
-          (double)Func::eval(val1_, val2_, val3_), (double)res_);
+      Kokkos::printf("value at %f, %f, %f which is %f was expected to be %f\n",
+                     (double)val1_, (double)val2_, (double)val3_,
+                     (double)Func::eval(val1_, val2_, val3_), (double)res_);
     }
   }
 };
@@ -1083,37 +1080,36 @@ struct TestAbsoluteValueFunction {
     using Kokkos::abs;
     if (abs(1) != 1 || abs(-1) != 1) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(int)\n");
+      Kokkos::printf("failed abs(int)\n");
     }
     if (abs(2l) != 2l || abs(-2l) != 2l) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(long int)\n");
+      Kokkos::printf("failed abs(long int)\n");
     }
     if (abs(3ll) != 3ll || abs(-3ll) != 3ll) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(long long int)\n");
+      Kokkos::printf("failed abs(long long int)\n");
     }
     if (abs(4.f) != 4.f || abs(-4.f) != 4.f) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(float)\n");
+      Kokkos::printf("failed abs(float)\n");
     }
     if (abs(5.) != 5. || abs(-5.) != 5.) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(double)\n");
+      Kokkos::printf("failed abs(double)\n");
     }
 #ifdef MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
     if (abs(6.l) != 6.l || abs(-6.l) != 6.l) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(long double)\n");
+      Kokkos::printf("failed abs(long double)\n");
     }
 #endif
     // special values
     using Kokkos::isinf;
     using Kokkos::isnan;
     if (abs(-0.) != 0. || !isinf(abs(-INFINITY)) || !isnan(abs(-NAN))) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "failed abs(floating_point) special values\n");
+      Kokkos::printf("failed abs(floating_point) special values\n");
     }
 
     static_assert(std::is_same<decltype(abs(1)), int>::value, "");
@@ -1145,14 +1141,14 @@ struct TestIsNaN {
     using Kokkos::Experimental::signaling_NaN;
     if (isnan(1) || isnan(INT_MAX)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(integral)\n");
+      Kokkos::printf("failed isnan(integral)\n");
     }
     if (isnan(2.f) || !isnan(quiet_NaN<float>::value) ||
         !isnan(signaling_NaN<float>::value)
 
     ) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(float)\n");
+      Kokkos::printf("failed isnan(float)\n");
     }
     if (isnan(3.)
 #ifndef KOKKOS_COMPILER_NVHPC  // FIXME_NVHPC
@@ -1161,20 +1157,19 @@ struct TestIsNaN {
 #endif
     ) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(double)\n");
+      Kokkos::printf("failed isnan(double)\n");
     }
 #ifdef MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
     if (isnan(4.l) || !isnan(quiet_NaN<long double>::value) ||
         !isnan(signaling_NaN<long double>::value)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(long double)\n");
+      Kokkos::printf("failed isnan(long double)\n");
     }
 #endif
     // special values
     if (isnan(INFINITY) || !isnan(NAN)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "failed isnan(floating_point) special values\n");
+      Kokkos::printf("failed isnan(floating_point) special values\n");
     }
 
     static_assert(std::is_same<decltype(isnan(1)), bool>::value, "");

diff --git a/core/unit_test/TestPrintf.hpp b/core/unit_test/TestPrintf.hpp
@@ -0,0 +1,33 @@
+//@HEADER
+// ************************************************************************
+//
+//                        Kokkos v. 4.0
+//       Copyright (2022) National Technology & Engineering
+//               Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#include <gtest/gtest.h>
+
+#include <Kokkos_Core.hpp>
+
+template <class ExecutionSpace>
+void test_kokkos_printf() {
+  ::testing::internal::CaptureStdout();
+  Kokkos::parallel_for(
+      Kokkos::RangePolicy<ExecutionSpace>(0, 1),
+      KOKKOS_LAMBDA(int) { Kokkos::printf("Print an integer: %d", 2); });
+  Kokkos::fence();
+  auto const captured = ::testing::internal::GetCapturedStdout();
+  std::string expected_string("Print an integer: 2");
+  ASSERT_EQ(captured, expected_string);
+}
+
+TEST(TEST_CATEGORY, kokkos_printf) { test_kokkos_printf<TEST_EXECSPACE>(); }
diff --git a/core/unit_test/TestRange.hpp b/core/unit_test/TestRange.hpp
@@ -134,8 +134,7 @@ struct TestRange {
   KOKKOS_INLINE_FUNCTION
   void operator()(const VerifyInitTag &, const int i) const {
     if (i != m_flags(i)) {
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("TestRange::test_for_error at %d != %d\n",
-                                    i, m_flags(i));
+      Kokkos::printf("TestRange::test_for_error at %d != %d\n", i, m_flags(i));
     }
   }
 
@@ -147,8 +146,7 @@ struct TestRange {
   KOKKOS_INLINE_FUNCTION
   void operator()(const VerifyResetTag &, const int i) const {
     if (2 * i != m_flags(i)) {
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("TestRange::test_for_error at %d != %d\n",
-                                    i, m_flags(i));
+      Kokkos::printf("TestRange::test_for_error at %d != %d\n", i, m_flags(i));
     }
   }
 
@@ -160,8 +158,8 @@ struct TestRange {
   KOKKOS_INLINE_FUNCTION
   void operator()(const VerifyOffsetTag &, const int i) const {
     if (i + offset != m_flags(i)) {
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("TestRange::test_for_error at %d != %d\n",
-                                    i + offset, m_flags(i));
+      Kokkos::printf("TestRange::test_for_error at %d != %d\n", i + offset,
+                     m_flags(i));
     }
   }