SWDEV-289409 - Add Infra for process isolation (#2374)

Change-Id: Iada6e5cfd38e2ba6efa14d7328e56f2260a72931

catch/README.md

@@ -80,6 +80,15 @@ Tests fall in 5 categories and its file name prefix are as follows:
  - Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment
  - Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc  y): Performance tests are used to get results of HIP APIs.
 
+There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn of process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf tests.
+Sample Usage:
+```cpp
+hip::SpawnProc proc(<relative path of exe with test folder>, <optional bool value, if output is to be recorded>);
+REQUIRE(0 == proc.run()); // Test of return value of the proc
+REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process
+```
+The process can be a standalone exe (see tests/catch/unit/printfExe for more information).
+
 General Guidelines:
  - Do not use the catch2 tags. Tags wont be used for filtering
  - Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails)

catch/hipTestMain/CMakeLists.txt

@@ -36,6 +36,7 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         OccupancyTest
                                         DeviceTest
                                         RTC
+                                        printfTests
                                         TextureTest
                                         stdc++fs)
 

catch/hipTestMain/hip_test_context.cc

@@ -105,7 +105,7 @@ bool TestContext::skipTest() const {
   return false;
 }
 
-std::string TestContext::currentPath() { return fs::current_path().string(); }
+std::string TestContext::currentPath() const { return fs::current_path().string(); }
 
 bool TestContext::parseJsonFile() {
   // Check if file exists

catch/include/hip_test_checkers.hh

@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2021 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.
+*/
+
 #pragma once
 #include "hip_test_common.hh"
 #include <iostream>

catch/include/hip_test_context.hh

@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2021 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.
+*/
+
 #pragma once
 #include <hip/hip_runtime.h>
 #include <vector>
@@ -72,7 +94,7 @@ class TestContext {
   bool skipTest() const;
 
   const std::string& getCurrentTest() const { return current_test; }
-  std::string currentPath();
+  std::string currentPath() const;
 
   TestContext(const TestContext&) = delete;
   void operator=(const TestContext&) = delete;

catch/include/hip_test_process.hh

@@ -0,0 +1,94 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+#pragma once
+#include "hip_test_common.hh"
+
+#include <string>
+#include <array>
+#include <cstdlib>
+#include <random>
+#include <fstream>
+#include <streambuf>
+
+#if __has_include(<filesystem>)
+#include <filesystem>
+namespace fs = std::filesystem;
+#elif __has_include(<experimental/filesystem>)
+#include <experimental/filesystem>
+namespace fs = std::experimental::filesystem;
+#else
+#error "gg filesystem"
+#endif
+
+namespace hip {
+class SpawnProc {
+  std::string exeName;
+  std::string resultStr;
+  std::string tmpFileName;
+  bool captureOutput;
+
+  std::string getRandomString(size_t len = 6) {
+    std::random_device dev;
+    std::mt19937 rng(dev());
+    std::uniform_int_distribution<std::mt19937::result_type> dist(0, 25);
+
+    std::string res;
+    for (size_t i = 0; i < len; i++) {
+      res += 'a' + dist(rng);
+    }
+    return res;
+  }
+
+ public:
+  SpawnProc(std::string exeName_, bool captureOutput_ = false)
+      : exeName(exeName_), captureOutput(captureOutput_) {
+    auto dir = fs::path(TestContext::get().currentPath()).parent_path();
+    dir /= exeName;
+    exeName = dir.string();
+    if (captureOutput) {
+      auto path = fs::temp_directory_path();
+      path /= getRandomString();
+      tmpFileName = path.string();
+    }
+  }
+
+  int run() {
+    std::string execCmd = exeName;
+    if (captureOutput) {
+      execCmd += " > ";
+      execCmd += tmpFileName;
+    }
+
+    auto res = std::system(execCmd.c_str());
+
+    if (captureOutput) {
+      std::ifstream t(tmpFileName.c_str());
+      resultStr =
+          std::string((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
+    }
+    return res;
+  }
+
+  std::string getOutput() { return resultStr; }
+};
+}  // namespace hip

catch/unit/CMakeLists.txt

@@ -25,4 +25,6 @@ add_subdirectory(event)
 add_subdirectory(occupancy)
 add_subdirectory(device)
 add_subdirectory(rtc)
+add_subdirectory(printf)
+add_subdirectory(printfExe)
 add_subdirectory(texture)

catch/unit/printf/CMakeLists.txt

@@ -0,0 +1,16 @@
+# AMD Tests
+set(TEST_SRC
+    printfFlags.cc
+    printfSpecifiers.cc
+)
+
+# Create shared lib of all tests
+add_library(printfTests SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+if(HIP_PLATFORM MATCHES "amd")
+    set_property(TARGET printfTests PROPERTY CXX_STANDARD 17)
+else()
+    target_compile_options(printfTests PUBLIC -std=c++17)
+endif()
+
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests printfTests)

catch/unit/printf/printfFlags.cc

@@ -0,0 +1,57 @@
+/*
+Copyright (c) 2021 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>
+
+__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);
+}
+
+TEST_CASE("Unit_printf_flags") {
+  std::string reference(R"here(00000042
+-0000042
+00000042
+0123.456
++0000042
+-42
++0000042
+xyzzy   
+-42
+00000042        
+        00000042
+)here");
+
+  hip::SpawnProc proc("unit/printfExe/printfFlags", true);
+  REQUIRE(proc.run() == 0);
+  REQUIRE(proc.getOutput() == reference);
+}

catch/unit/printf/printfSpecifiers.cc

@@ -0,0 +1,95 @@
+/*
+Copyright (c) 2021 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>
+
+TEST_CASE("Unit_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("unit/printfExe/printfSepcifiers", true);
+  REQUIRE(0 == proc.run());
+  REQUIRE(proc.getOutput() == reference);
+}

catch/unit/printfExe/CMakeLists.txt

@@ -0,0 +1,5 @@
+add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags.cc)
+add_executable(printfSepcifiers EXCLUDE_FROM_ALL printfSepcifiers.cc)
+
+add_dependencies(build_tests printfFlags)
+add_dependencies(build_tests printfSepcifiers)

catch/unit/printfExe/printfFlags.cc

@@ -0,0 +1,42 @@
+/*
+Copyright (c) 2021 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>>>();
+  hipDeviceSynchronize();
+}

catch/unit/printfExe/printfSepcifiers.cc

@@ -0,0 +1,65 @@
+/*
+Copyright (c) 2021 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>>>();
+  hipDeviceSynchronize();
+  return 0;
+}