AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add HIP_CHECK_THREAD and REQUIRE_THREAD macro for multi threaded HIP API tests (#2664)

Browse Source
This commit is contained in:
Jatin Chaudhary
2022-06-20 10:37:13 +01:00
committed by GitHub
parent 902eae64e3
commit 5afcd13390
7 changed files with 414 additions and 101 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/catch/include/hip_test_checkers.hh
+141 -43
View File
@@ -23,17 +23,17 @@ THE SOFTWARE.
#pragma once
#include "hip_test_common.hh"
#include <iostream>
#include<fstream>
#include<regex>
#include <fstream>
#include <regex>
#include <type_traits>
#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<typename T> // pointer type
bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) {
template <typename T> // 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 <typename T> 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<T, char>::value || std::is_same<T, unsigned char>::value) {
} else if (std::is_same<T, char>::value || std::is_same<T, unsigned char>::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 <typename T> 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<T, int>::value || std::is_same<T, unsigned int>::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<T, char>::value || std::is_same<T, unsigned char>::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 <typename T>
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 <typename T>
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 <typename T> 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 <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (usePinnedHost) {
if (A_h) {
@@ -210,6 +314,21 @@ template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin
return true;
}
template <typename T>
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 <typename T>
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 <typename T>
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<typename T>
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<T, float>::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 {
tests/catch/include/hip_test_common.hh
+91 -42
View File
@@ -22,9 +22,14 @@ THE SOFTWARE.
#pragma once
#include "hip_test_context.hh"
#include <hip_test_rtc.hh>
#include <catch.hpp>
#include <atomic>
#include <chrono>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <mutex>
#include <cstdlib>
#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 <iostream>
#include <iomanip>
#include <chrono>
#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 <typename... Typenames, typename K, typename Dim, typename... Args>
void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock,
hipStream_t stream, Args&&... packedArgs) {
#ifndef RTC_TESTING
validateArguments(kernel, packedArgs...);
kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
validateArguments(kernel, packedArgs...);
kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
#else
launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
std::forward<Args>(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; \
}
tests/catch/include/hip_test_context.hh
+28
View File
@@ -23,7 +23,11 @@ THE SOFTWARE.
#pragma once
#include <hip/hip_runtime.h>
#include <hip/hiprtc.h>
#include <atomic>
#include <mutex>
#include <vector>
#include <iostream>
#include <string>
#include <set>
#include <unordered_map>
@@ -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<HCResult> results; // Multi threaded test results buffer
std::atomic<bool> 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();
};