EXSWCPHIPT-42: Changes to HIP RTC Framework implementation (#2732)

- Removed ifdef from hipTestContext class
- Fix potential race condition in hipTest::launchRTCKernel()
- Improve documentation
- Move moduleUnloading to main() instead of explicitly calling it on every test
- Fix code formating
- Fix segmentation fault caused by using catch2 macro after catch2 is destroyed

[ROCm/hip commit: d46d399976]

projects/hip/tests/README.md

@@ -149,19 +149,25 @@ Here "-C performance" indicate the "performance" configuration of ctest.
 
 ### RTC Testing
 
-To enable RTC testing, cmake needs to be passed the DRTC_TESTING=1 options.
+To enable RTC testing, cmake needs to be passed the `-DRTC_TESTING=1` option.
 
 When this option is passed, all tests that support this functionality will be run using HIP RTC to compile and run.
 
 To enable HIP RTC support for a specific test:
-    1 - Move all its kernels to tests/catch/kernels (one file per kernel)
-    2 - Update tests/catch/kernels/CMakeLists.txt
-    3 - Update tests/catch/include/kernels.hh
-    4 - Update tests/catch/include/kernel_mapping.hh
-    5 - Include kernels.hh
-    6 - Call hipTest::launchKernel() function instead of hipLaunchKernelGGL()
 
-Note: HIP RTC does not do implicit casting of kernel parameters. This requires the test writer to explicitly do all the casting before running the kernel. The code will not compile otherwise.
+1. Move all its kernels to `tests/catch/kernels` (one file per kernel):
+    1. Kernel **functions** should use the file extension `.cpp` and include `kernels.hh`
+    2. Kernel **templates** should use the file extension `.inl`
+2. Update `tests/catch/kernels/CMakeLists.txt` (i.e. add the new kernel **functions** to `TEST_SRC`)
+3. Update `tests/catch/include/kernels.hh`:
+    1. Declare the new kernel **functions**
+    2. Include the new .inl files that contain kernel **templates**
+    3. Call the `FUNCTION_WRAPPER` and `TEMPLATE_WRAPPER` macros for each new function and template respectively.
+4. Update `tests/catch/include/kernel_mapping.hh` with the mapping between the new files and respective function / template names.
+5. Include `kernels.hh`
+6. Call the `hipTest::launchKernel()` function instead of `hipLaunchKernelGGL()`
+
+**Note:** HIP RTC does not do implicit casting of kernel parameters. This **requires** the test writer to explicitly do all the casting before running the kernel. There is a `static_assert` inside `hipTest::launchKernel()` that checks that this was done correctly. However, due to limitations, the assertion is only performed when  `-DRTC_TESTING` option is **disabled**. This means that runtime errors can occur if the casts are not performed correctly and `-DRTC_TESTING` is enabled.
 
 ### If a test fails - how to debug a test
 

projects/hip/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake

@@ -204,9 +204,9 @@ function(hip_add_exe_to_target)
   )
   # Create shared lib of all tests
   if(NOT RTC_TESTING)
-  add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object> $<TARGET_OBJECTS:KERNELS>)
+    add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object> $<TARGET_OBJECTS:KERNELS>)
   else ()
-  add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object>)
+    add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object>)
     if(HIP_PLATFORM STREQUAL "amd")
       target_link_libraries(${_NAME} hiprtc)
     else()

projects/hip/tests/catch/hipTestMain/hip_test_context.cc

@@ -1,3 +1,4 @@
+#include <cstdlib>
 #include <hip_test_common.hh>
 #include <picojson.h>
 #include <fstream>
@@ -211,7 +212,10 @@ bool TestContext::parseJsonFile() {
 
 void TestContext::cleanContext() {
   for (auto& pair : compiledKernels) {
-    REQUIRE(hipSuccess == hipModuleUnload(pair.second.module));
+    hipError_t error = hipModuleUnload(pair.second.module);
+    if (error != hipSuccess) {
+      throw std::runtime_error("Unable to unload rtc module");
+    }
   }
 }
 

projects/hip/tests/catch/hipTestMain/main.cc

@@ -9,5 +9,8 @@ int main(int argc, char** argv) {
     std::cout << "HIP_SKIP_THIS_TEST" << std::endl;
     return 0;
   }
-  return Catch::Session().run(argc, argv);
+  int out = Catch::Session().run(argc, argv);
+  TestContext::get().cleanContext();
+  return out;
+  
 }

projects/hip/tests/catch/include/hip_test_rtc.hh

@@ -31,6 +31,7 @@ THE SOFTWARE.
 #include <fstream>
 #include <sstream>
 #include <set>
+#include <mutex>
 #include "hip/hip_runtime_api.h"
 #include "hip_test_context.hh"
 
@@ -218,7 +219,6 @@ static inline void printInfo() {
 template <typename... Typenames, typename... Args>
 void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThreads,
                      std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) {
-
   printInfo();
   TestContext& testContext = TestContext::get();
   std::string kernelName = (*getKernelName)();
@@ -226,25 +226,29 @@ void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThr
   std::vector<std::string> kernelTypenames{std::string(HipTest::getTypeName<Typenames>())...};
   std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames);
 
-  if (testContext.getFunction(kernelExpression) == nullptr) {
-    hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)};
-    std::vector<char> compiledCode{getKernelCode(rtcProgram)};
+  static std::mutex mutex{};
+  {
+    std::lock_guard<std::mutex> lockGuard(mutex);
+    if (testContext.getFunction(kernelExpression) == nullptr) {
+      hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)};
+      std::vector<char> compiledCode{getKernelCode(rtcProgram)};
 
-    hipModule_t module;
+      hipModule_t module;
 
-    REQUIRE(hipSuccess == hipModuleLoadData(&module, compiledCode.data()));
+      REQUIRE(hipSuccess == hipModuleLoadData(&module, compiledCode.data()));
 
       hipFunction_t kernelFunction;
 
-    const char* loweredName;
-    REQUIRE(HIPRTC_SUCCESS ==
-          hiprtcGetLoweredName(rtcProgram, kernelExpression.c_str(), &loweredName));
-    REQUIRE(hipSuccess == hipModuleGetFunction(&kernelFunction, module, loweredName));
+      const char* loweredName;
+      REQUIRE(HIPRTC_SUCCESS ==
+              hiprtcGetLoweredName(rtcProgram, kernelExpression.c_str(), &loweredName));
+      REQUIRE(hipSuccess == hipModuleGetFunction(&kernelFunction, module, loweredName));
 
-    /* After obtaining the kernelFunction, the program is no longer needed. So it can be destroyed */
-    REQUIRE(HIPRTC_SUCCESS == hiprtcDestroyProgram(&rtcProgram));
+      /* After obtaining the kernelFunction, the program is no longer needed. So it can be destroyed */
+      REQUIRE(HIPRTC_SUCCESS == hiprtcDestroyProgram(&rtcProgram));
 
-    testContext.trackRtcState(kernelExpression, module, kernelFunction);
+      testContext.trackRtcState(kernelExpression, module, kernelFunction);
+    }
   }
 
   hipFunction_t kernelFunction = testContext.getFunction(kernelExpression);

projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc

@@ -87,5 +87,4 @@ static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
     HIP_CHECK(hipEventDestroy(stop));
 
     HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
-    TestContext::get().cleanContext();
 }

projects/hip/tests/catch/unit/memory/hipHostMalloc.cc

@@ -144,7 +144,6 @@ TEST_CASE("Unit_hipHostMalloc_Basic") {
     HIP_CHECK(hipHostFree(A_h));
     HIP_CHECK(hipHostFree(B_h));
     HIP_CHECK(hipHostFree(C_h));
-    TestContext::get().cleanContext();
   }
 }
 /*
@@ -182,7 +181,6 @@ TEST_CASE("Unit_hipHostMalloc_NonCoherent") {
                    SYNC_STREAM, ptrType);
   CheckHostPointer(numElements, A, hipEventReleaseToSystem,
                    SYNC_EVENT, ptrType);
-  TestContext::get().cleanContext();
 }
 
 /*