diff --git a/tests/catch/README.md b/tests/catch/README.md index ade70de8e9..7d1ff6265b 100644 --- a/tests/catch/README.md +++ b/tests/catch/README.md @@ -34,6 +34,22 @@ Some useful functions are: This information can be accessed in any test via using: `TestContext::get().isAmd()`. +## Adding test for a specific platform +There might be some functionality which is not present on some platforms. Those tests can be hidden inside following macros. + +- ```HT_AMD``` is 1 when tests are running on AMD platform and 0 on NVIDIA. +- ```HT_NVIDIA``` is 1 when tests are running on NVIDIA platform and 0 on AMD + +Usage: + +```cpp +#if HT_AMD +TEST_CASE("hipExtAPIs") { + // ... +} +#endif +``` + ## Config file schema Some tests can be skipped using a config file placed in hipTestMain/config folder. Multiple config files can be defined for different configurations. The naming convention for the file needs to be "config_platform_os_archname.json" @@ -56,16 +72,108 @@ The schema of the json file is as follows: } ``` -## Env Variables +## Environment Variables - `HT_CONFIG_FILE` : This variable can be set to the config file name or full path. Disabled tests will be read from this. - `HT_LOG_ENABLE` : This is for debugging the HIP Test Framework itself. Setting it to 1, all `LogPrintf` will be printed on screen +## Test Macros +### Single Thread Macros +These macros are to be used when your test is calling HIP APIs via the main thread. + +- `HIP_CHECK` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. + + - Usage: ```HIP_CHECK(hipMalloc(&dPtr, 10));``` + +- ```HIP_CHECK_ERROR``` : This macro takes in a HIP API and tests its result against a provided result. This can be used when the API is expected to fail with a particular result. + + - Usage: ```HIP_CHECK_ERROR(hipMalloc(&dPtr, 0), hipErrorInvalidValue);``` + +- ```HIPRTC_CHECK``` : This macro takes in a HIPRTC API and tests its result against HIPRTC_SUCCESS. + + - Usage: ```HIPRTC_CHECK(hiprtcCompileProgram(prog, count, options));``` + +- ```HIP_ASSERT``` : This macro takes in a bool condition as input and does a ```REQUIRE``` on the condition. + + - Usage: ```HIP_ASSERT(result == 10);``` + +### Multi Thread Macros +These macros are to be used when you call HIP APIs in a multi threaded way. They exist because Catch2 ```REQUIRE``` and ```CHECK``` macros can not handle multi threaded calls. To solve this problem, two macros are added```HIP_CHECK_THREAD``` and ```REQUIRE_THREAD``` which can be used to check result of HIP APIs and test assertions respectively. The results can be validate after the threads join via ```HIP_CHECK_THREAD_FINALIZE```. + +Note: These should used in ```std::thread``` only. For multi proc guidelines look at [MultiProc Macros](#multi-process-macros) and [SpawnProc Class](#multiproc-management-class) + +- ```HIP_CHECK_THREAD``` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. It can also tell other threads if an error has occured in one of the HIP API and can prematurely stop the threads. + +- ```REQUIRE_THREAD``` : This macro takes in a bool condition and tests for its result to be true. If this check fails, it can signal other threads to terminate early. + +- ```HIP_CHECK_THREAD_FINALIZE``` : This macro checks for the results logged by ```HIP_CHECK_THREAD```. This needs to be called after the threads have joined. + +Please also note that you can not return values in functions calling ```HIP_CHECK_THREAD``` or ```REQUIRE_THREAD``` macro. + + Usage: + + ```cpp + auto threadFunc = []() { + int *dPtr{nullptr}; + HIP_CHECK_THREAD(hipMalloc(&dPtr, 10)); + REQUIRE_THREAD(dPtr != nullptr); + // Some other work + }; + + // Launch threads + std::vector threadPool; + for(...) { + threadPool.emplace_back(std::thread(threadFunc)); + } + + // Join threads + for(auto &i : threadPool) { + i.join(); + } + + // Validate all results + HIP_CHECK_THREAD_FINALIZE(); + ``` + +### Skipping Tests if certain criteria is not met +If there arises a condition where certain flag is disabled and due to which a test can not run at that time, the following macro can be of use. It will highlight the test in ctest report as well. + +- ```HIP_SKIP_TEST``` : The api takes in an input of the reason as well and prints out the line HIP_SKIP_THIS_TEST. This causes ctest to mark the test as skipped and the test shows up in the report as skipped prompting proper response from the team. + + Usage: + + ```cpp + TEST_CASE("TestOnlyOnXnack") { + if(!XNACKEnabled) { + HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); + return; + } + // Rest of test functionality + } + ``` + +### Multi Process Macros +These macros are to be called in multi process tests, inside a process which gets spawned. The reasoning is the same, Catch2 does not support multi process checks. + +- ```HIPCHECK``` : Same as ```HIP_CHECK``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +- ```HIPASSERT``` : Same as ```HIP_ASSERT``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +## MultiProc Management Class +There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output. +Sample Usage: +```cpp +hip::SpawnProc proc(, ); +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). + ## Enabling New Tests -Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=ON```. After porting existing tests, this will be turned on by default. +Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default. ## Building a single test ```bash -hipcc -I/tests/newTests/include /tests/newTests/hipTestMain/standalone_main.cc -I/tests/newTests/external/Catch2 -g -o +hipcc -I/tests/catch/include /tests/catch/hipTestMain/standalone_main.cc -I/tests/catch/external/Catch2 -g -o ``` ## Debugging support @@ -87,16 +195,7 @@ 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(, ); -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: +# 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) - Check return of each HIP API and fail whenever there is a misma tch with hipSuccess or hiprtcSuccess. diff --git a/tests/catch/hipTestMain/hip_test_context.cc b/tests/catch/hipTestMain/hip_test_context.cc index a6e3a08609..597813b4c8 100644 --- a/tests/catch/hipTestMain/hip_test_context.cc +++ b/tests/catch/hipTestMain/hip_test_context.cc @@ -229,4 +229,42 @@ hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) { } else { return nullptr; } +} + +void TestContext::addResults(HCResult r) { + std::unique_lock lock(resultMutex); + results.push_back(r); + if ((!r.conditionsResult) || + ((r.result != hipSuccess) && (r.result != hipErrorPeerAccessAlreadyEnabled))) { + hasErrorOccured_.store(true); + } +} + +void TestContext::finalizeResults() { + std::unique_lock lock(resultMutex); + // clear the results whatever happens + std::shared_ptr emptyVec(nullptr, [this](auto) { results.clear(); }); + + for (const auto& i : results) { + INFO("HIP API Result check\n File:: " + << i.file << "\n Line:: " << i.line << "\n API:: " << i.call + << "\n Result:: " << i.result << "\n Result Str:: " << hipGetErrorString(i.result)); + REQUIRE(((i.result == hipSuccess) || (i.result == hipErrorPeerAccessAlreadyEnabled))); + REQUIRE(i.conditionsResult); + } + hasErrorOccured_.store(false); // Clear the flag +} + +bool TestContext::hasErrorOccured() { return hasErrorOccured_.load(); } + +TestContext::~TestContext() { + // Show this message when there are unchecked results + if (results.size() != 0) { + std::cerr << "HIP_CHECK_THREAD_FINALIZE() has not been called after HIP_CHECK_THREAD\n" + << "Please call HIP_CHECK_THREAD_FINALIZE after joining threads\n" + << "There is/are " << results.size() << " unchecked results from threads." + << std::endl; + std::abort(); // Crash to bring users attention to this message and avoid accidental passing of + // tests without checking for errors + } } \ No newline at end of file diff --git a/tests/catch/include/hip_test_checkers.hh b/tests/catch/include/hip_test_checkers.hh index 77ac0f4d74..fce61fe677 100644 --- a/tests/catch/include/hip_test_checkers.hh +++ b/tests/catch/include/hip_test_checkers.hh @@ -23,17 +23,17 @@ THE SOFTWARE. #pragma once #include "hip_test_common.hh" #include -#include -#include +#include +#include #include -#define guarantee(cond, str) \ - { \ - if (!(cond)) { \ - INFO("guarantee failed: " << str); \ - abort(); \ - } \ - } +#define guarantee(cond, str) \ + { \ + if (!(cond)) { \ + INFO("guarantee failed: " << str); \ + abort(); \ + } \ + } namespace HipTest { @@ -73,15 +73,15 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM return mismatchCount; } -template // pointer type -bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) { +template // pointer type +bool checkArray(T* hData, T* hOutputData, size_t width, size_t height, size_t depth = 1) { for (size_t i = 0; i < depth; i++) { for (size_t j = 0; j < height; j++) { for (size_t k = 0; k < width; k++) { - int offset = i*width*height + j*width + k; + int offset = i * width * height + j * width + k; if (hData[offset] != hOutputData[offset]) { - INFO("Mismatch at [" << i << "," << j << "," << k << "]:" - << hData[offset] << "----" << hOutputData[offset]); + INFO("Mismatch at [" << i << "," << j << "," << k << "]:" << hData[offset] << "----" + << hOutputData[offset]); CHECK(false); return false; } @@ -120,7 +120,7 @@ template void setDefaultData(size_t numElements, T* A_h, T* B_h, T* if (A_h) A_h[i] = 3; if (B_h) B_h[i] = 4; if (C_h) C_h[i] = 5; - } else if(std::is_same::value || std::is_same::value) { + } else if (std::is_same::value || std::is_same::value) { if (A_h) A_h[i] = 'a'; if (B_h) B_h[i] = 'b'; if (C_h) C_h[i] = 'c'; @@ -185,6 +185,110 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N, return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost); } +// Threaded version of setDefaultData to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template void setDefaultDataT(size_t numElements, T* A_h, T* B_h, T* C_h) { + // Initialize the host data: + + for (size_t i = 0; i < numElements; i++) { + if (std::is_same::value || std::is_same::value) { + if (A_h) A_h[i] = 3; + if (B_h) B_h[i] = 4; + if (C_h) C_h[i] = 5; + } else if (std::is_same::value || std::is_same::value) { + if (A_h) A_h[i] = 'a'; + if (B_h) B_h[i] = 'b'; + if (C_h) C_h[i] = 'c'; + } else { + if (A_h) A_h[i] = 3.146f + i; + if (B_h) B_h[i] = 1.618f + i; + if (C_h) C_h[i] = 1.4f + i; + } + } +} + +// Threaded version of initArraysForHost to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template +void initArraysForHostT(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (usePinnedHost) { + if (A_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)A_h, Nbytes)); + } + if (B_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)B_h, Nbytes)); + } + if (C_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)C_h, Nbytes)); + } + } else { + if (A_h) { + *A_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*A_h != nullptr); + } + + if (B_h) { + *B_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*B_h != nullptr); + } + + if (C_h) { + *C_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*C_h != nullptr); + } + } + + setDefaultDataT(N, A_h ? *A_h : nullptr, B_h ? *B_h : nullptr, C_h ? *C_h : nullptr); +} + +// Threaded version of initArrays to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template +void initArraysT(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N, + bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (A_d) { + HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes)); + } + if (B_d) { + HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes)); + } + if (C_d) { + HIP_CHECK_THREAD(hipMalloc(C_d, Nbytes)); + } + + initArraysForHostT(A_h, B_h, C_h, N, usePinnedHost); +} + +// Threaded version of freeArraysForHost to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template void freeArraysForHostT(T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (usePinnedHost) { + if (A_h) { + HIP_CHECK_THREAD(hipHostFree(A_h)); + } + if (B_h) { + HIP_CHECK_THREAD(hipHostFree(B_h)); + } + if (C_h) { + HIP_CHECK_THREAD(hipHostFree(C_h)); + } + } else { + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + if (C_h) { + free(C_h); + } + } +} + template bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) { if (usePinnedHost) { if (A_h) { @@ -210,6 +314,21 @@ template bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin return true; } +template +void freeArraysT(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (A_d) { + HIP_CHECK_THREAD(hipFree(A_d)); + } + if (B_d) { + HIP_CHECK_THREAD(hipFree(B_d)); + } + if (C_d) { + HIP_CHECK_THREAD(hipFree(C_d)); + } + + freeArraysForHostT(A_h, B_h, C_h, usePinnedHost); +} + template bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) { if (A_d) { @@ -226,20 +345,6 @@ bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo } template -unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock, - size_t N) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device)); - - unsigned blocks = props.multiProcessorCount * blocksPerCU; - if (blocks * threadsPerBlock > N) { - blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - } - return blocks; -} -template static bool assemblyFile_Verification(std::string assemfilename, std::string inst) { std::string filePath = "./catch/unit/deviceLib/"; bool result = false; @@ -254,34 +359,27 @@ static bool assemblyFile_Verification(std::string assemfilename, std::string ins while (getline(file, line)) { line_pos++; if ((std::is_same::value)) { - if (!start_pos && - std::regex_search(line, - std::regex("Begin function (.*)AtomicCheck"))) { + if (!start_pos && std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) { start_pos = line_pos; } - if (!last_pos && - std::regex_search(line, - std::regex(".Lfunc_end0-(.*)AtomicCheck"))) { + if (!last_pos && std::regex_search(line, std::regex(".Lfunc_end0-(.*)AtomicCheck"))) { last_pos = line_pos; break; } } else { - if ((start_match != 2) && std::regex_search(line, - std::regex("Begin function (.*)AtomicCheck"))) { + if ((start_match != 2) && + std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) { start_match++; - if (start_match == 2) - start_pos = line_pos; + if (start_match == 2) start_pos = line_pos; } - if (!last_pos && std::regex_search(line, - std::regex("func_end1-(.*)AtomicCheck"))) { + if (!last_pos && std::regex_search(line, std::regex("func_end1-(.*)AtomicCheck"))) { last_pos = line_pos; break; } } if (start_pos) { result = std::regex_search(line, std::regex(inst)); - if (result) - break; + if (result) break; } } } else { diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 63b8354421..fe272f7ace 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -22,9 +22,14 @@ THE SOFTWARE. #pragma once #include "hip_test_context.hh" -#include #include +#include +#include #include +#include +#include +#include +#include #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); @@ -33,22 +38,51 @@ THE SOFTWARE. { \ hipError_t localError = error; \ if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ - INFO("Error: " << hipGetErrorString(localError) << " Code: " << localError << " Str: " \ - << #error << " In File: " << __FILE__ << " At line: " << __LINE__); \ + INFO("Error: " << hipGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ REQUIRE(false); \ } \ } +// Threaded HIP_CHECKs +#define HIP_CHECK_THREAD(error) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localError = error; \ + HCResult result(__LINE__, __FILE__, localError, #error); \ + TestContext::get().addResults(result); \ + } + +#define REQUIRE_THREAD(condition) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localResult = (condition); \ + HCResult result(__LINE__, __FILE__, hipSuccess, #condition, localResult); \ + TestContext::get().addResults(result); \ + } + +// Do not call before all threads have joined +#define HIP_CHECK_THREAD_FINALIZE() \ + { TestContext::get().finalizeResults(); } + + // Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. #define HIP_CHECK_ERROR(errorExpr, expectedError) \ { \ hipError_t localError = errorExpr; \ INFO("Matching Errors: " \ - << " Expected Error: " << hipGetErrorString(expectedError) \ - << " Expected Code: " << expectedError << '\n' \ + << "\n Expected Error: " << hipGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ << " Actual Error: " << hipGetErrorString(localError) \ - << " Actual Code: " << localError << "\nStr: " << #errorExpr \ - << "\nIn File: " << __FILE__ << " At line: " << __LINE__); \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ REQUIRE(localError == expectedError); \ } @@ -57,8 +91,9 @@ THE SOFTWARE. { \ auto localError = error; \ if (localError != HIPRTC_SUCCESS) { \ - INFO("Error: " << hiprtcGetErrorString(localError) << " Code: " << localError << " Str: " \ - << #error << " In File: " << __FILE__ << " At line: " << __LINE__); \ + INFO("Error: " << hiprtcGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ REQUIRE(false); \ } \ } @@ -67,12 +102,6 @@ THE SOFTWARE. #define HIP_ASSERT(x) \ { REQUIRE((x)); } -#ifdef __cplusplus -#include -#include -#include -#endif - #define HIPCHECK(error) \ { \ hipError_t localError = error; \ @@ -84,19 +113,20 @@ THE SOFTWARE. } #define HIPASSERT(condition) \ - if (!(condition)) { \ - printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ - abort(); \ - } + if (!(condition)) { \ + printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ + abort(); \ + } + #if HT_NVIDIA -#define CTX_CREATE() \ - hipCtx_t context;\ +#define CTX_CREATE() \ + hipCtx_t context; \ initHipCtx(&context); #define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context)); #define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array)); #define HIP_TEX_REFERENCE hipTexRef #define HIP_ARRAY hiparray -static void initHipCtx(hipCtx_t *pcontext) { +static void initHipCtx(hipCtx_t* pcontext) { HIPCHECK(hipInit(0)); hipDevice_t device; HIPCHECK(hipDeviceGet(&device, 0)); @@ -130,9 +160,9 @@ static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) { } static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) { - int device; + int device{0}; HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; + hipDeviceProp_t props{}; HIP_CHECK(hipGetDeviceProperties(&props, device)); unsigned blocks = props.multiProcessorCount * blocksPerCU; @@ -143,23 +173,40 @@ static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlo return blocks; } -static inline int RAND_R(unsigned* rand_seed) -{ - #if defined(_WIN32) || defined(_WIN64) - srand(*rand_seed); - return rand(); - #else - return rand_r(rand_seed); - #endif +// Threaded version of setNumBlocks - to be used in multi threaded test +// Why? because catch2 does not support multithreaded macro calls +// Make sure you call HIP_CHECK_THREAD_FINALIZE after your threads join +// Also you can not return in threaded functions, due to how HIP_CHECK_THREAD works +static inline void setNumBlocksThread(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N, + unsigned& blocks) { + int device{0}; + blocks = 0; // incase error has occured in some other thread and the next call might not execute, + // we set the blocks size to 0 + HIP_CHECK_THREAD(hipGetDevice(&device)); + hipDeviceProp_t props{}; + HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device)); + + blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } +} + +static inline int RAND_R(unsigned* rand_seed) { +#if defined(_WIN32) || defined(_WIN64) + srand(*rand_seed); + return rand(); +#else + return rand_r(rand_seed); +#endif } inline bool isImageSupported() { - int imageSupport = 1; + int imageSupport = 1; #ifdef __HIP_PLATFORM_AMD__ - int device; - HIP_CHECK(hipGetDevice(&device)); - HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - device)); + int device; + HIP_CHECK(hipGetDevice(&device)); + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device)); #endif return imageSupport != 0; } @@ -217,8 +264,8 @@ template void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) { #ifndef RTC_TESTING - validateArguments(kernel, packedArgs...); - kernel<<>>(std::forward(packedArgs)...); + validateArguments(kernel, packedArgs...); + kernel<<>>(std::forward(packedArgs)...); #else launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, std::forward(packedArgs)...); @@ -229,6 +276,8 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB // This must be called in the beginning of image test app's main() to indicate whether image // is supported. -#define checkImageSupport() \ - if (!HipTest::isImageSupported()) \ - { printf("Texture is not support on the device. Skipped.\n"); return; } +#define checkImageSupport() \ + if (!HipTest::isImageSupported()) { \ + printf("Texture is not support on the device. Skipped.\n"); \ + return; \ + } diff --git a/tests/catch/include/hip_test_context.hh b/tests/catch/include/hip_test_context.hh index 2d630b70f7..adcaf25aad 100644 --- a/tests/catch/include/hip_test_context.hh +++ b/tests/catch/include/hip_test_context.hh @@ -23,7 +23,11 @@ THE SOFTWARE. #pragma once #include #include + +#include +#include #include +#include #include #include #include @@ -64,6 +68,18 @@ typedef struct Config_ { std::string os; // windows/linux } Config; +// Store Multi threaded results +struct HCResult { + size_t line; // Line of check (HIP_CHECK_THREAD or REQUIRE_THREAD) + std::string file; // File name of the check + hipError_t result; // hipResult for HIP_CHECK_THREAD, for conditions its hipSuccess + std::string call; // Call of HIP API or a bool condition + bool conditionsResult; // If bool condition, result of call. For HIP Calls its true + HCResult(size_t l, std::string f, hipError_t r, std::string c, bool b = true) + : line(l), file(f), result(r), call(c), conditionsResult(b) {} +}; + + class TestContext { bool p_windows = false, p_linux = false; // OS bool amd = false, nvidia = false; // HIP Platform @@ -97,6 +113,11 @@ class TestContext { TestContext(int argc, char** argv); + // Multi threaded checks helpers + std::mutex resultMutex; + std::vector results; // Multi threaded test results buffer + std::atomic hasErrorOccured_{false}; + public: static TestContext& get(int argc = 0, char** argv = nullptr) { static TestContext instance(argc, argv); @@ -112,6 +133,11 @@ class TestContext { const std::string& getCurrentTest() const { return current_test; } std::string currentPath() const; + // Multi threaded results helpers + void addResults(HCResult r); // Add multi threaded results + void finalizeResults(); // Validate on all results + bool hasErrorOccured(); // Query if error has occured + /** * @brief Unload all loaded modules. * Note: This function needs to be called at the end of each test that uses RTC. @@ -142,4 +168,6 @@ class TestContext { TestContext(const TestContext&) = delete; void operator=(const TestContext&) = delete; + + ~TestContext(); }; diff --git a/tests/catch/unit/stream/streamCommon.cc b/tests/catch/unit/stream/streamCommon.cc index 265142e23b..14ac4000eb 100644 --- a/tests/catch/unit/stream/streamCommon.cc +++ b/tests/catch/unit/stream/streamCommon.cc @@ -95,11 +95,11 @@ __global__ void waiting_kernel(int* semaphore) { std::thread startSignalingThread(int* semaphore) { std::thread signalingThread([semaphore]() { hipStream_t signalingStream; - HIP_CHECK(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); + HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore); - HIP_CHECK(hipStreamSynchronize(signalingStream)); - HIP_CHECK(hipStreamDestroy(signalingStream)); + HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream)); + HIP_CHECK_THREAD(hipStreamDestroy(signalingStream)); }); return signalingThread; diff --git a/tests/catch/unit/stream/streamCommon.hh b/tests/catch/unit/stream/streamCommon.hh index db73f1f668..1d5a1ea958 100644 --- a/tests/catch/unit/stream/streamCommon.hh +++ b/tests/catch/unit/stream/streamCommon.hh @@ -44,6 +44,7 @@ __global__ void waiting_kernel(int* semaphore = nullptr); /** * @brief Creates a thread that runs a signaling_kernel on a non-blocking stream. * hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream. + * You must call HIP_CHECK_THREAD_FINALIZE after joining this thread. * * @param semaphore memory location to signal * @return std::thread thread that has to be joined after the testing is done.