Remove KOKKOS_IMPL_DO_NOT_USE_PRINTF (kokkos#6593)

* Remove KOKKOS_IMPL_DO_NOT_USE_PRINTF * Clean up tutorials
masterleinad · Nov 17, 2023 · 81a9586 · 81a9586
1 parent ff7104c
commit 81a9586
Show file tree

Hide file tree

Showing 9 changed files with 47 additions and 85 deletions.
diff --git a/core/src/Kokkos_Macros.hpp b/core/src/Kokkos_Macros.hpp
@@ -339,12 +339,6 @@
 #define KOKKOS_IMPL_DEVICE_FUNCTION
 #endif
 
-// 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__)
-#endif
-
 //----------------------------------------------------------------------------
 // Define final version of functions. This is so that clang tidy can find these
 // macros more easily

diff --git a/core/src/setup/Kokkos_Setup_SYCL.hpp b/core/src/setup/Kokkos_Setup_SYCL.hpp
@@ -38,12 +38,4 @@
 #include <CL/sycl.hpp>
 #endif
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define KOKKOS_IMPL_DO_NOT_USE_PRINTF(format, ...)                \
-  do {                                                            \
-    const __attribute__((opencl_constant)) char fmt[] = (format); \
-    sycl::ext::oneapi::experimental::printf(fmt, ##__VA_ARGS__);  \
-  } while (0)
-#endif
-
 #endif
diff --git a/core/unit_test/TestBitManipulationBuiltins.hpp b/core/unit_test/TestBitManipulationBuiltins.hpp
@@ -804,26 +804,26 @@ struct TestBitCastFunction {
     using Kokkos::bit_cast;
     if (bit_cast<int>(123) != 123) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #1\n");
+      Kokkos::printf("failed check #1\n");
     }
     if (bit_cast<int>(123u) != 123) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #2\n");
+      Kokkos::printf("failed check #2\n");
     }
     if (bit_cast<int>(~0u) != ~0) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #3\n");
+      Kokkos::printf("failed check #3\n");
     }
     if constexpr (sizeof(int) == sizeof(float)) {
       if (!check<int>(12.34f)) {
         ++e;
-        KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #4\n");
+        Kokkos::printf("failed check #4\n");
       }
     }
     if constexpr (sizeof(unsigned long long) == sizeof(double)) {
       if (!check<unsigned long long>(123.456)) {
         ++e;
-        KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #5\n");
+        Kokkos::printf("failed check #5\n");
       }
     }
 
@@ -848,11 +848,11 @@ struct TestBitCastFunction {
     }
     if (!(bit_cast<S>(arr) == arr)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #6\n");
+      Kokkos::printf("failed check #6\n");
     }
     if (!(bit_cast<S>(arr2) == arr2)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed check #7\n");
+      Kokkos::printf("failed check #7\n");
     }
   }
 };

diff --git a/core/unit_test/TestMathematicalFunctions.hpp b/core/unit_test/TestMathematicalFunctions.hpp
@@ -1304,12 +1304,12 @@ struct TestAbsoluteValueFunction {
     if (abs(static_cast<KE::half_t>(4.f)) != static_cast<KE::half_t>(4.f) ||
         abs(static_cast<KE::half_t>(-4.f)) != static_cast<KE::half_t>(4.f)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(KE::half_t)\n");
+      Kokkos::printf("failed abs(KE::half_t)\n");
     }
     if (abs(static_cast<KE::bhalf_t>(4.f)) != static_cast<KE::bhalf_t>(4.f) ||
         abs(static_cast<KE::bhalf_t>(-4.f)) != static_cast<KE::bhalf_t>(4.f)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed abs(KE::bhalf_t)\n");
+      Kokkos::printf("failed abs(KE::bhalf_t)\n");
     }
     if (abs(5.) != 5. || abs(-5.) != 5.) {
       ++e;
@@ -1360,35 +1360,34 @@ struct TestFloatingPointAbsoluteValueFunction {
     using Kokkos::fabs;
     if (fabs(4.f) != 4.f || fabs(-4.f) != 4.f) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fabs(float)\n");
+      Kokkos::printf("failed fabs(float)\n");
     }
     if (fabs(static_cast<KE::half_t>(4.f)) != static_cast<KE::half_t>(4.f) ||
         fabs(static_cast<KE::half_t>(-4.f)) != static_cast<KE::half_t>(4.f)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fabs(KE::half_t)\n");
+      Kokkos::printf("failed fabs(KE::half_t)\n");
     }
     if (fabs(static_cast<KE::bhalf_t>(4.f)) != static_cast<KE::bhalf_t>(4.f) ||
         fabs(static_cast<KE::bhalf_t>(-4.f)) != static_cast<KE::bhalf_t>(4.f)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fabs(KE::bhalf_t)\n");
+      Kokkos::printf("failed fabs(KE::bhalf_t)\n");
     }
     if (fabs(5.) != 5. || fabs(-5.) != 5.) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fabs(double)\n");
+      Kokkos::printf("failed fabs(double)\n");
     }
 #ifdef MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
     if (fabs(6.l) != 6.l || fabs(-6.l) != 6.l) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fabs(long double)\n");
+      Kokkos::printf("failed fabs(long double)\n");
     }
 #endif
     // special values
     using Kokkos::isinf;
     using Kokkos::isnan;
     if (fabs(-0.) != 0. || !isinf(fabs(-INFINITY)) || !isnan(fabs(-NAN))) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "failed fabs(floating_point) special values\n");
+      Kokkos::printf("failed fabs(floating_point) special values\n");
     }
 
     static_assert(std::is_same<decltype(fabs(static_cast<KE::half_t>(4.f))),
@@ -1420,7 +1419,7 @@ struct TestFloatingPointRemainderFunction : FloatingPointComparison {
     if (!compare(fmod(6.2f, 4.f), 2.2f, 1) &&
         !compare(fmod(-6.2f, 4.f), -2.2f, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fmod(float)\n");
+      Kokkos::printf("failed fmod(float)\n");
     }
     if (!compare(
             fmod(static_cast<KE::half_t>(6.2f), static_cast<KE::half_t>(4.f)),
@@ -1429,7 +1428,7 @@ struct TestFloatingPointRemainderFunction : FloatingPointComparison {
             fmod(static_cast<KE::half_t>(-6.2f), static_cast<KE::half_t>(4.f)),
             -static_cast<KE::half_t>(2.2f), 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fmod(KE::half_t)\n");
+      Kokkos::printf("failed fmod(KE::half_t)\n");
     }
     if (!compare(
             fmod(static_cast<KE::bhalf_t>(6.2f), static_cast<KE::bhalf_t>(4.f)),
@@ -1438,17 +1437,17 @@ struct TestFloatingPointRemainderFunction : FloatingPointComparison {
                       static_cast<KE::bhalf_t>(4.f)),
                  -static_cast<KE::bhalf_t>(2.2f), 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fmod(KE::bhalf_t)\n");
+      Kokkos::printf("failed fmod(KE::bhalf_t)\n");
     }
     if (!compare(fmod(6.2, 4.), 2.2, 1) && !compare(fmod(-6.2, 4.), -2.2, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fmod(double)\n");
+      Kokkos::printf("failed fmod(double)\n");
     }
 #ifdef MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
     if (!compare(fmod(6.2l, 4.l), 2.2l, 1) &&
         !compare(fmod(-6.2l, 4.l), -2.2l, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed fmod(long double)\n");
+      Kokkos::printf("failed fmod(long double)\n");
     }
 #endif
     // special values
@@ -1457,8 +1456,7 @@ struct TestFloatingPointRemainderFunction : FloatingPointComparison {
     if (!isinf(fmod(-KE::infinity<float>::value, 1.f)) &&
         !isnan(fmod(-KE::quiet_NaN<float>::value, 1.f))) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
-          "failed fmod(floating_point) special values\n");
+      Kokkos::printf("failed fmod(floating_point) special values\n");
     }
 
     static_assert(std::is_same<decltype(fmod(static_cast<KE::half_t>(4.f),
@@ -1494,7 +1492,7 @@ struct TestIEEEFloatingPointRemainderFunction : FloatingPointComparison {
     if (!compare(remainder(6.2f, 4.f), 2.2f, 2) &&
         !compare(remainder(-6.2f, 4.f), 2.2f, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed remainder(float)\n");
+      Kokkos::printf("failed remainder(float)\n");
     }
     if (!compare(remainder(static_cast<KE::half_t>(6.2f),
                            static_cast<KE::half_t>(4.f)),
@@ -1503,7 +1501,7 @@ struct TestIEEEFloatingPointRemainderFunction : FloatingPointComparison {
                            static_cast<KE::half_t>(4.f)),
                  -static_cast<KE::half_t>(2.2f), 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed remainder(KE::half_t)\n");
+      Kokkos::printf("failed remainder(KE::half_t)\n");
     }
     if (!compare(remainder(static_cast<KE::bhalf_t>(6.2f),
                            static_cast<KE::bhalf_t>(4.f)),
@@ -1512,18 +1510,18 @@ struct TestIEEEFloatingPointRemainderFunction : FloatingPointComparison {
                            static_cast<KE::bhalf_t>(4.f)),
                  -static_cast<KE::bhalf_t>(2.2f), 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed remainder(KE::bhalf_t)\n");
+      Kokkos::printf("failed remainder(KE::bhalf_t)\n");
     }
     if (!compare(remainder(6.2, 4.), 2.2, 2) &&
         !compare(remainder(-6.2, 4.), 2.2, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed remainder(double)\n");
+      Kokkos::printf("failed remainder(double)\n");
     }
 #ifdef MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
     if (!compare(remainder(6.2l, 4.l), 2.2l, 1) &&
         !compare(remainder(-6.2l, 4.l), -2.2l, 1)) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed remainder(long double)\n");
+      Kokkos::printf("failed remainder(long double)\n");
     }
 #endif
     // special values
@@ -1532,7 +1530,7 @@ struct TestIEEEFloatingPointRemainderFunction : FloatingPointComparison {
     if (!isinf(remainder(-KE::infinity<float>::value, 1.f)) &&
         !isnan(remainder(-KE::quiet_NaN<float>::value, 1.f))) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF(
+      Kokkos::printf(
           "failed remainder(floating_point) special values\n");
     }
 
@@ -1748,7 +1746,7 @@ struct TestIsNaN {
 #endif
     ) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(KE::half_t)\n");
+      Kokkos::printf("failed isnan(KE::half_t)\n");
     }
     if (isnan(static_cast<KE::bhalf_t>(2.f))
 #ifndef KOKKOS_COMPILER_NVHPC  // FIXME_NVHPC 23.7
@@ -1758,7 +1756,7 @@ struct TestIsNaN {
 #endif
     ) {
       ++e;
-      KOKKOS_IMPL_DO_NOT_USE_PRINTF("failed isnan(KE::bhalf_t)\n");
+      Kokkos::printf("failed isnan(KE::bhalf_t)\n");
     }
     if (isnan(3.)
 #ifndef KOKKOS_COMPILER_NVHPC  // FIXME_NVHPC 23.7

diff --git a/example/tutorial/01_hello_world/hello_world.cpp b/example/tutorial/01_hello_world/hello_world.cpp
@@ -58,12 +58,7 @@ struct hello_world {
   // is unnecessary but harmless.
   KOKKOS_INLINE_FUNCTION
   void operator()(const int i) const {
-    // FIXME_SYCL needs workaround for printf
-#ifndef __SYCL_DEVICE_ONLY__
-    printf("Hello from i = %i\n", i);
-#else
-    (void)i;
-#endif
+    Kokkos::printf("Hello from i = %i\n", i);
   }
 };
 

diff --git a/example/tutorial/01_hello_world_lambda/hello_world_lambda.cpp b/example/tutorial/01_hello_world_lambda/hello_world_lambda.cpp
@@ -76,13 +76,9 @@ int main(int argc, char* argv[]) {
 #if defined(KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA)
   Kokkos::parallel_for(
       15, KOKKOS_LAMBDA(const int i) {
-  // FIXME_SYCL needs workaround for printf
-#ifndef __SYCL_DEVICE_ONLY__
-        // printf works in a CUDA parallel kernel; std::ostream does not.
-        printf("Hello from i = %i\n", i);
-#else
-	(void)i;
-#endif
+        // Kokko::printf works for all backends in a parallel kernel;
+        // std::ostream does not.
+        Kokkos::printf("Hello from i = %i\n", i);
       });
 #endif
   // You must call finalize() after you are done using Kokkos.

diff --git a/example/tutorial/Hierarchical_Parallelism/01_thread_teams/thread_teams.cpp b/example/tutorial/Hierarchical_Parallelism/01_thread_teams/thread_teams.cpp
@@ -47,13 +47,9 @@ struct hello_world {
     // The TeamPolicy<>::member_type provides functions to query the multi
     // dimensional index of a thread as well as the number of thread-teams and
     // the size of each team.
-#ifndef __SYCL_DEVICE_ONLY__
-    // FIXME_SYCL needs printf workaround
-    printf("Hello World: %i %i // %i %i\n", thread.league_rank(),
-           thread.team_rank(), thread.league_size(), thread.team_size());
-#else
-    (void)thread;
-#endif
+    Kokkos::printf("Hello World: %i %i // %i %i\n", thread.league_rank(),
+                   thread.team_rank(), thread.league_size(),
+                   thread.team_size());
   }
 };
 

diff --git a/example/tutorial/Hierarchical_Parallelism/01_thread_teams_lambda/thread_teams_lambda.cpp b/example/tutorial/Hierarchical_Parallelism/01_thread_teams_lambda/thread_teams_lambda.cpp
@@ -57,16 +57,12 @@ int main(int narg, char* args[]) {
       policy,
       KOKKOS_LAMBDA(const team_member& thread, int& lsum) {
         lsum += 1;
-    // TeamPolicy<>::member_type provides functions to query the
-    // multidimensional index of a thread, as well as the number of
-    // thread teams and the size of each team.
-#ifndef __SYCL_DEVICE_ONLY__
-        // FIXME_SYCL needs workaround for printf
-        printf("Hello World: %i %i // %i %i\n", thread.league_rank(),
-               thread.team_rank(), thread.league_size(), thread.team_size());
-#else
-        (void)thread;
-#endif
+        // TeamPolicy<>::member_type provides functions to query the
+        // multidimensional index of a thread, as well as the number of
+        // thread teams and the size of each team.
+        Kokkos::printf("Hello World: %i %i // %i %i\n", thread.league_rank(),
+                       thread.team_rank(), thread.league_size(),
+                       thread.team_size());
       },
       sum);
 #endif

diff --git a/example/tutorial/Hierarchical_Parallelism/02_nested_parallel_for/nested_parallel_for.cpp b/example/tutorial/Hierarchical_Parallelism/02_nested_parallel_for/nested_parallel_for.cpp
@@ -43,16 +43,11 @@ struct hello_world {
     // the operator using a team_policy acts like a parallel region for the
     // team. That means that everything outside of the nested parallel_for is
     // also executed by all threads of the team.
-    Kokkos::parallel_for(Kokkos::TeamThreadRange(thread, 31),
-                         [&](const int& i) {
-#ifndef __SYCL_DEVICE_ONLY__
-                           // FIXME_SYCL needs printf workaround
-                           printf("Hello World: (%i , %i) executed loop %i \n",
-                                  thread.league_rank(), thread.team_rank(), i);
-#else
-			   (void) i;
-#endif
-                         });
+    Kokkos::parallel_for(
+        Kokkos::TeamThreadRange(thread, 31), [&](const int& i) {
+          Kokkos::printf("Hello World: (%i , %i) executed loop %i \n",
+                         thread.league_rank(), thread.team_rank(), i);
+        });
   }
 };