From cc7ae079015ca04abc87c01f9e3de66d1d1c4e3a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Tue, 28 Jun 2022 04:30:37 +0100 Subject: [PATCH] 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: d46d399976c173521c9af0247fc80c71221ebf74] --- projects/hip/tests/README.md | 22 +++++++++----- .../external/Catch2/cmake/Catch2/Catch.cmake | 4 +-- .../catch/hipTestMain/hip_test_context.cc | 6 +++- projects/hip/tests/catch/hipTestMain/main.cc | 5 +++- .../hip/tests/catch/include/hip_test_rtc.hh | 30 +++++++++++-------- .../catch/unit/event/Unit_hipEventRecord.cc | 1 - .../tests/catch/unit/memory/hipHostMalloc.cc | 2 -- 7 files changed, 42 insertions(+), 28 deletions(-) diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md index 1b1e54a2e3..e03e399104 100644 --- a/projects/hip/tests/README.md +++ b/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 diff --git a/projects/hip/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake b/projects/hip/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake index 603dc1b994..868ccfa739 100644 --- a/projects/hip/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/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} $ $) + add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $ $) else () - add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $) + add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $) if(HIP_PLATFORM STREQUAL "amd") target_link_libraries(${_NAME} hiprtc) else() diff --git a/projects/hip/tests/catch/hipTestMain/hip_test_context.cc b/projects/hip/tests/catch/hipTestMain/hip_test_context.cc index 597813b4c8..ae496d372b 100644 --- a/projects/hip/tests/catch/hipTestMain/hip_test_context.cc +++ b/projects/hip/tests/catch/hipTestMain/hip_test_context.cc @@ -1,3 +1,4 @@ +#include #include #include #include @@ -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"); + } } } diff --git a/projects/hip/tests/catch/hipTestMain/main.cc b/projects/hip/tests/catch/hipTestMain/main.cc index 886aa7a8dc..ea3bb9bd23 100644 --- a/projects/hip/tests/catch/hipTestMain/main.cc +++ b/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; + } diff --git a/projects/hip/tests/catch/include/hip_test_rtc.hh b/projects/hip/tests/catch/include/hip_test_rtc.hh index 0860ae51db..0f862840bb 100644 --- a/projects/hip/tests/catch/include/hip_test_rtc.hh +++ b/projects/hip/tests/catch/include/hip_test_rtc.hh @@ -31,6 +31,7 @@ THE SOFTWARE. #include #include #include +#include #include "hip/hip_runtime_api.h" #include "hip_test_context.hh" @@ -218,7 +219,6 @@ static inline void printInfo() { template 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 kernelTypenames{std::string(HipTest::getTypeName())...}; std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames); - if (testContext.getFunction(kernelExpression) == nullptr) { - hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)}; - std::vector compiledCode{getKernelCode(rtcProgram)}; + static std::mutex mutex{}; + { + std::lock_guard lockGuard(mutex); + if (testContext.getFunction(kernelExpression) == nullptr) { + hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)}; + std::vector 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); diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc b/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc index d2f13d6e12..408f802d41 100644 --- a/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc +++ b/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -87,5 +87,4 @@ static_cast(A_d), static_cast(B_d), C_d, N); HIP_CHECK(hipEventDestroy(stop)); HipTest::checkVectorADD(A_h, B_h, C_h, N, true); - TestContext::get().cleanContext(); } diff --git a/projects/hip/tests/catch/unit/memory/hipHostMalloc.cc b/projects/hip/tests/catch/unit/memory/hipHostMalloc.cc index 39a2bb6719..f0b1fafe8d 100644 --- a/projects/hip/tests/catch/unit/memory/hipHostMalloc.cc +++ b/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(); } /*