SWDEV-403479 - [catch2][dtest] Adding Buffered compiler option to Exi…

…sting Printf Tests

Change-Id: I4fc422bd7d5ee1896743978fe078ac3f8d57e319

catch/include/hip_test_defgroups.hh

@@ -142,3 +142,10 @@ THE SOFTWARE.
  * This section describes tests for the Vector type functions and operators.
  * @}
  */
+
+/**
+ * @defgroup PrintfTest Printf API Management
+ * @{
+ * This section describes the various Printf use case Scenarios.
+ * @}
+ */

catch/unit/printf/CMakeLists.txt

@@ -2,8 +2,15 @@
 set(TEST_SRC
     printfFlags.cc
     printfSpecifiers.cc
+    printfFlagsNonHost.cc
+    printfSpecifiersNonHost.cc
 )
-
+if(UNIX)
+  if(HIP_PLATFORM MATCHES "amd")
+    set_source_files_properties(printfFlagsNonHost.cc PROPERTIES COMPILE_OPTIONS "-mprintf-kind=buffered")
+    set_source_files_properties(printfSpecifiersNonHost.cc PROPERTIES COMPILE_OPTIONS "-mprintf-kind=buffered")
+  endif()
+endif()
 
 if(HIP_PLATFORM MATCHES "amd")
   hip_add_exe_to_target(NAME printfTests
@@ -20,6 +27,11 @@ endif()
 # Standalone exes
 add_executable(printfFlags_exe EXCLUDE_FROM_ALL printfFlags_exe.cc)
 add_executable(printfSpecifiers_exe EXCLUDE_FROM_ALL printfSpecifiers_exe.cc)
+add_executable(printfFlagsNonHost_exe EXCLUDE_FROM_ALL printfFlagsNonHost_exe.cc)
+add_executable(printfSpecifiersNonHost_exe EXCLUDE_FROM_ALL printfSpecifiersNonHost_exe.cc)
 
 add_dependencies(build_tests printfFlags_exe)
 add_dependencies(build_tests printfSpecifiers_exe)
+add_dependencies(build_tests printfFlagsNonHost_exe)
+add_dependencies(build_tests printfSpecifiersNonHost_exe)
+

catch/unit/printf/printfFlagsNonHost.cc

@@ -0,0 +1,61 @@
+/*Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_process.hh>
+#include <hip_test_defgroups.hh>
+
+/**
+* @addtogroup printf
+* @{
+* @ingroup PrintfTest
+* `int printf()` -
+* Method to print the content on output device.
+*/
+/**
+* Test Description
+* ------------------------
+* - Test case to verify the printf return value from other process for the compiler option -mprintf-kind=buffered
+* - Fetch the printf content from a process. Compare it with reference string.
+* Test source
+* ------------------------
+* - catch/unit/printf/printfFlagsNonHost.cc
+* Test requirements
+* ------------------------
+* - HIP_VERSION >= 5.7
+*/
+
+TEST_CASE("Unit_Buffered_Printf_Flags") {
+  std::string reference(R"here(00000042
+-0000042
+00000042
+0123.456
++0000042
+-42
++0000042
+xyzzy   
+-42
+00000042        
+        00000042
+)here");
+
+  hip::SpawnProc proc("printfFlagsNonHost_exe", true);
+  REQUIRE(proc.run() == 0);
+  REQUIRE(proc.getOutput() == reference);
+}
+

catch/unit/printf/printfFlagsNonHost_exe.cc

@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip/hip_runtime.h>
+
+__global__ void test_kernel() {
+  printf("%08d\n", 42);
+  printf("%08i\n", -42);
+  printf("%08u\n", 42);
+  printf("%08g\n", 123.456);
+  printf("%0+8d\n", 42);
+  printf("%+d\n", -42);
+  printf("%+08d\n", 42);
+  printf("%-8s\n", "xyzzy");
+  printf("% i\n", -42);
+  printf("%-16.8d\n", 42);
+  printf("%16.8d\n", 42);
+}
+
+int main() {
+  test_kernel<<<1, 1>>>();
+  static_cast<void>(hipDeviceSynchronize());
+}

catch/unit/printf/printfSpecifiersNonHost.cc

@@ -0,0 +1,116 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_process.hh>
+#include <hip_test_defgroups.hh>
+
+/**
+* @addtogroup printf
+* @{
+* @ingroup PrintfTest
+* `int printf()` -
+* Method to print the content on output device.
+*/
+/**
+* Test Description
+* ------------------------
+* - Test case to verify the different format specifiers. Test case should compile with the compiler option -mprintf-kind=buffered
+* - Fetch the printf content from a process which will verify format specifier. Compare it with reference string.
+* Test source
+* ------------------------
+* - catch/unit/printf/printfSpecifiersNonHost.cc
+* Test requirements
+* ------------------------
+* - HIP_VERSION >= 5.7
+*/
+
+TEST_CASE("Unit_Buffered_Printf_Specifier") {
+#ifdef __HIP_PLATFORM_NVIDIA__
+  std::string reference(R"here(xyzzy
+%
+hello % world
+%s
+%s0xf01dab1eca55e77e
+%cxyzzy
+sep
+-42
+42
+123.456000
+-123.456000
+-1.234560e+02
+1.234560E+02
+123.456
+-123.456
+x
+(null)
+(nil)
+3.14159000    hello 0xf01dab1eca55e77e
+)here");
+#elif !defined(_WIN32)
+  std::string reference(R"here(xyzzy
+%
+hello % world
+%s
+%s0xf01dab1eca55e77e
+%cxyzzy
+sep
+-42
+42
+123.456000
+-123.456000
+-1.234560e+02
+1.234560E+02
+123.456
+-123.456
+x
+
+(nil)
+3.14159000    hello 0xf01dab1eca55e77e
+)here");
+#else
+  std::string reference(R"here(xyzzy
+%
+hello % world
+%s
+%sF01DAB1ECA55E77E
+%cxyzzy
+sep
+-42
+42
+123.456000
+-123.456000
+-1.234560e+02
+1.234560E+02
+123.456
+-123.456
+x
+
+0000000000000000
+3.14159000    hello F01DAB1ECA55E77E
+)here");
+#endif
+
+  hip::SpawnProc proc("printfSpecifiersNonHost_exe", true);
+  REQUIRE(0 == proc.run());
+  REQUIRE(proc.getOutput() == reference);
+}

catch/unit/printf/printfSpecifiersNonHost_exe.cc

@@ -0,0 +1,66 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip/hip_runtime.h>
+
+__global__ void test_kernel() {
+  const char* N = nullptr;
+  const char* s = "hello world";
+  printf("xyzzy\n");
+  printf("%%\n");
+  printf("hello %% world\n");
+  printf("%%s\n");
+  // Two special tests to make sure that the compiler pass correctly
+  // skips over a '%%' without affecting the logic for locating
+  // string arguments.
+  printf("%%s%p\n", (void*)0xf01dab1eca55e77e);
+  printf("%%c%s\n", "xyzzy");
+  printf("%c%c%c\n", 's', 'e', 'p');
+  printf("%d\n", -42);
+  printf("%u\n", 42);
+  printf("%f\n", 123.456);
+#ifdef __HIP_PLATFORM_AMD__
+  printf("%F\n", -123.456);
+#else
+  printf("%f\n", -123.456);
+#endif
+  printf("%e\n", -123.456);
+  printf("%E\n", 123.456);
+  printf("%g\n", 123.456);
+  printf("%G\n", -123.456);
+  printf("%c\n", 'x');
+  printf("%s\n", N);
+  printf("%p\n", (void *)N);
+#ifdef __HIP_PLATFORM_AMD__
+  printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 5, s, (void*)0xf01dab1eca55e77e);
+#else
+  // In Cuda, printf doesn't support %.*, %*.*
+  printf("%.8f %8.5s %p\n", 3.14159, s, (void*)0xf01dab1eca55e77e);
+#endif
+}
+
+int main() {
+  test_kernel<<<1, 1>>>();
+  static_cast<void>(hipDeviceSynchronize());
+  return 0;
+}
+