diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 8518ff07cb..7d7c6de46f 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -43,4 +43,5 @@ endif() hip_add_exe_to_target(NAME StreamTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests) + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) diff --git a/catch/unit/stream/hipStreamCreateWithPriority.cc b/catch/unit/stream/hipStreamCreateWithPriority.cc index 300d78ac3c..9405db8289 100644 --- a/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -509,8 +509,8 @@ bool validateStreamPrioritiesWithEvents() { #define OP(x) \ free(src_h_##x); \ free(dst_h_##x); \ - hipFree(src_d_##x); \ - hipFree(dst_d_##x); + HIP_CHECK(hipFree(src_d_##x)); \ + HIP_CHECK(hipFree(dst_d_##x)); OP(low) OP(normal) OP(high) diff --git a/catch/unit/stream/hipStreamValue.cc b/catch/unit/stream/hipStreamValue.cc index 5eb4321fef..5e70122195 100644 --- a/catch/unit/stream/hipStreamValue.cc +++ b/catch/unit/stream/hipStreamValue.cc @@ -17,276 +17,377 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #include +#include +#include constexpr unsigned int writeFlag = 0; -#define DEFINE_HIP_STREAM_VALUE(TYPE, BITS, ...) hipStream##TYPE##Value##BITS(__VA_ARGS__) +template auto waitFunc(Args... args) { + if constexpr (std::is_same::value) { + return hipStreamWaitValue32(args...); + } else { + return hipStreamWaitValue64(args...); + } +}; -#define CHECK_HIP_STREAM_VALUE(TYPE, BITS, ...) \ - HIP_CHECK(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__)); +template auto writeFunc(Args... args) { + if constexpr (std::is_same::value) { + return hipStreamWriteValue32(args...); + } else { + return hipStreamWriteValue64(args...); + } +}; -#define NEG_TEST_ERROR_CHECK(TYPE, BITS, errorCode, ...) \ - HIP_CHECK_ERROR(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__), errorCode); +// Random predefined 32 and 64 bit values +using value32_t = std::integral_constant; +using value64_t = std::integral_constant; +template +using testValue = + typename std::conditional::value, value32_t, value64_t>::type; -#if HT_AMD -// Random predefiend 32 and 64 bit values -constexpr uint32_t value32 = 0x70F0F0FF; -constexpr uint64_t value64 = 0x7FFF0000FFFF0000; constexpr uint32_t DATA_INIT = 0x1234; constexpr uint32_t DATA_UPDATE = 0X4321; -template struct TEST_WAIT { - using uintT = typename std::make_unsigned::type; - int compareOp; - uintT mask; - uintT waitValue; - intT signalValueFail; - intT signalValuePass; +template struct TEST_WAIT { + static_assert(std::is_same::value or std::is_same::value, + "only implemented for 32 bit and 64 bit unsigned integers"); + unsigned int compareOp; + UIntT mask = ~static_cast(0); + UIntT waitValue; + UIntT signalValueFail; + UIntT signalValuePass; - TEST_WAIT(int compareOp, uintT waitValue, intT signalValueFail, intT signalValuePass) + TEST_WAIT(unsigned int compareOp, UIntT waitValue, UIntT signalValueFail, UIntT signalValuePass) : compareOp{compareOp}, waitValue{waitValue}, signalValueFail{signalValueFail}, - signalValuePass{signalValuePass} { - mask = static_cast(0xFFFFFFFFFFFFFFFF); - } + signalValuePass{signalValuePass} {} - TEST_WAIT(int compareOp, uintT mask, uintT waitValue, intT signalValueFail, intT signalValuePass) + TEST_WAIT(unsigned int compareOp, UIntT mask, UIntT waitValue, UIntT signalValueFail, + UIntT signalValuePass) : compareOp{compareOp}, mask{mask}, waitValue{waitValue}, signalValueFail{signalValueFail}, signalValuePass{signalValuePass} {} }; -typedef TEST_WAIT TEST_WAIT32; -typedef TEST_WAIT TEST_WAIT64; + +using TEST_WAIT32 = TEST_WAIT; +using TEST_WAIT64 = TEST_WAIT; bool streamWaitValueSupported() { int device_num = 0; HIP_CHECK(hipGetDeviceCount(&device_num)); - int waitValueSupport; for (int device_id = 0; device_id < device_num; ++device_id) { HIP_CHECK(hipSetDevice(device_id)); - waitValueSupport = 0; - HIP_CHECK(hipDeviceGetAttribute(&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, - device_id)); + int waitValueSupport = 0; + auto getAttributeError = hipDeviceGetAttribute( + &waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, device_id); + if (getAttributeError != hipSuccess) { + HipTest::HIP_SKIP_TEST("attribute not supported"); + return false; + } if (waitValueSupport == 1) return true; } return false; } -// hipStreamWriteValue Tests -TEST_CASE("Unit_hipStreamValue_Write") { - int64_t* signalPtr; +// The different types of memory that can be used with hipStream[Wait|Write] +enum class PtrType { HostPtr, DevicePtr, DevicePtrToHost, Signal }; - hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); +// Helper class to expose the pointer that is used with hipStream[Write|Wait]Value and also store a +// unique pointer with the deleter to simplify cleanup +// Also includes functions to update and get the value directly +template class TestPtr { + // This stores the memory that must be deleted, as well as the deleter + UniquePtrWithDeleter ptrToDelete; - // Allocate Host Memory - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); + public: + // The pointer that should be used with hipStream[Write|Wait]Value + UIntT* ptr; - // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); + TestPtr(UIntT* ptr, UniquePtrWithDeleter ptrToDelete) + : ptrToDelete(std::move(ptrToDelete)), ptr(ptr) {} - // Register Signal Memory - HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory)); - // Initialise Data - *signalPtr = 0x0; - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; - - SECTION("Registered host memory hipStreamWriteValue32") { - INFO("Test writting to registered host pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue32(stream, hostPtr32.get(), value32, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr32 == value32); + // directly retrieve the value from wherever it was allocated + UIntT getValue(size_t offset = 0) { + if constexpr (type == PtrType::Signal || type == PtrType::HostPtr || + type == PtrType::DevicePtrToHost) { + return ptrToDelete.get()[offset]; + } else { + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + UIntT value; + HIP_CHECK(hipMemcpy(&value, ptr + offset, sizeof(UIntT), hipMemcpyDeviceToHost)); + return value; + } } - SECTION("Registered host memory hipStreamWriteValue64") { - INFO("Test writting to registered host pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue64(stream, hostPtr64.get(), value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr64 == value64); + // directly set the value wherever it was allocated + void setValue(UIntT value, size_t offset = 0) { + if constexpr (type == PtrType::Signal || type == PtrType::DevicePtrToHost || + type == PtrType::HostPtr) { + ptrToDelete.get()[offset] = value; + } else { + // hipMemcpy causes deadlock, so use hipStreamWriteValue + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(writeFunc(stream, ptr + offset, value, writeFlag)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } } +}; - // Test writting device pointer - void* devicePtr64; - void* devicePtr32; - HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr64, hostPtr64.get(), 0)); - HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr32, hostPtr32.get(), 0)); - // Reset values - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; +// required for the static assert +template inline constexpr bool AMD_ACTIVE = HT_AMD == 1; - SECTION("Device Memory hipStreamWriteValue32") { - INFO("Test writting to device pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue32(stream, devicePtr32, value32, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr32 == value32); +template auto allocMem() { + constexpr std::size_t arraySize = 1024; + if constexpr (type == PtrType::Signal) { + static_assert(std::is_same::value, + "signal memory should only be used with 64bit memory"); + + // Allocate Signal Memory + uint64_t* signalPtr{}; + + static_assert(AMD_ACTIVE, + "nvidia backend compiler doesn't like hipExtMallocWithFlags, even in this " + "constexpr branch"); +#if HT_AMD + // 8 is the only acceptable size + HIP_CHECK( + hipExtMallocWithFlags(reinterpret_cast(&signalPtr), 8, hipMallocSignalMemory)); +#endif + + // Init Memory + *signalPtr = 0; + + auto freeStuff = [](uint64_t* sPtr) { HIP_CHECK(hipFree(sPtr)); }; + return TestPtr>{ + signalPtr, std::unique_ptr(signalPtr, freeStuff)}; + } else if constexpr (type == PtrType::DevicePtrToHost) { + auto hostPtr = new UIntT[arraySize]; + + // Register Host Memory + HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0)); + + // Init memory + std::fill(hostPtr, hostPtr + arraySize, 0); + + UIntT* devicePtr; + // Test writing device pointer + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&devicePtr), hostPtr, 0)); + auto freeStuff = [](UIntT* ptr) { + HIP_CHECK(hipHostUnregister(ptr)); + delete[] ptr; + }; + + return TestPtr>{ + devicePtr, std::unique_ptr(hostPtr, freeStuff)}; + } else if constexpr (type == PtrType::HostPtr) { + auto hostPtr = new UIntT[arraySize]; + + // Register Host Memory + HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0)); + + // Init memory + std::fill(hostPtr, hostPtr + arraySize, 0); + + auto freeStuff = [](UIntT* ptr) { + HIP_CHECK(hipHostUnregister(ptr)); + delete[] ptr; + }; + + return TestPtr>{ + hostPtr, std::unique_ptr(hostPtr, freeStuff)}; + } else { + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + UIntT* devicePtr; + HIP_CHECK(hipMalloc(&devicePtr, sizeof(UIntT) * arraySize)); + HIP_CHECK(hipMemset(devicePtr, 0, sizeof(UIntT) * arraySize)); + auto freeStuff = [](UIntT* ptr) { HIP_CHECK(hipFree(ptr)); }; + return TestPtr>{ + devicePtr, std::unique_ptr(devicePtr, freeStuff)}; } - - SECTION("Device Memory hipStreamWriteValue64") { - INFO("Test writting to device pointer using hipStreamWriteValue64"); - HIP_CHECK(hipStreamWriteValue64(stream, devicePtr64, value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr64 == value64); - } - - // Test Writing to Signal Memory - SECTION("Signal Memory hipStreamWriteValue64") { - INFO("Test writting to signal memory using hipStreamWriteValue64"); - HIP_CHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*signalPtr == value64); - } - - // Cleanup - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipFree(signalPtr)); } -// hipStreamWaitValue Tests -template -void initData(intT* dataPtr, int64_t* signalPtr, TEST_T tc, std::vector& events) { - // Initialize memory to be waited on - *signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass; - - - // Initialize host pointers - dataPtr[0] = DATA_INIT; - dataPtr[1] = DATA_INIT; - - - hipEvent_t firstWriteEvent{nullptr}; - hipEvent_t secondWriteEvent{nullptr}; - HIP_CHECK(hipEventCreate(&firstWriteEvent)); - HIP_CHECK(hipEventCreate(&secondWriteEvent)); - events.push_back(firstWriteEvent); - events.push_back(secondWriteEvent); -} - -template -void syncAndCheckData(hipStream_t stream, intT* dataPtr, int64_t* signalPtr, TEST_T tc, - std::vector& events) { - // Ensure first part of host memory is updated - HIP_CHECK(hipStreamWaitEvent(stream, events[0], 0)); - HIP_ASSERT(dataPtr[0] == DATA_UPDATE); - if (isBlocking) { - // Ensure second part of host memory isn't updated yet - HIP_ASSERT(hipEventQuery(events[1]) == hipErrorNotReady); - HIP_ASSERT(dataPtr[1] == DATA_INIT); - // Update value to release stream - *signalPtr = tc.signalValuePass; +// allows the creation of a list of offsets while avoiding it for signal memory +template constexpr auto get_offsets() { + if constexpr (type == PtrType::Signal) { + return std::array{0}; + } else { + return std::array{0, 1, 2, 3, 31, 1023}; } - - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(hipEventQuery(events[1]) == hipSuccess); - // Finally ensure that second part of host memory is updated - HIP_ASSERT(dataPtr[1] == DATA_UPDATE); } -template void cleanup(hipStream_t& stream, intT* dataPtr, int64_t* signalPtr) { - // Cleanup - HIP_CHECK(hipFree(signalPtr)); - HIP_CHECK(hipHostUnregister(dataPtr)); - HIP_CHECK(hipStreamDestroy(stream)); -} +template struct TestParams { + using UIntType = UIntT; + constexpr static PtrType ptrType = ptrTypeValue; +}; -template void testWait(TEST_T tc) { +#if HT_AMD +TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams)) { +#else +TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams)) { +#endif +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-126"); + return; +#endif if (!streamWaitValueSupported()) { - UNSCOPED_INFO(" hipStreamWaitValue: not supported on this device , skipping ..."); + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); return; } - // Initialize stream + using UIntT = typename TestType::UIntType; + constexpr auto ptrType = TestType::ptrType; + constexpr auto writeValue = testValue::value; + hipStream_t stream{nullptr}; HIP_CHECK(hipStreamCreate(&stream)); + const auto offsets = get_offsets(); + const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets))); + + CAPTURE(offset); + // Allocate Host Memory - std::unique_ptr dataPtr(new intT(2)); + auto ptr = allocMem(); + UIntT* target = ptr.ptr + offset; + HIP_CHECK(writeFunc(stream, target, writeValue, writeFlag)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(ptr.getValue(offset) == writeValue); - // Register Host Memory - HIP_CHECK(hipHostRegister(&(dataPtr.get()[0]), sizeof(intT), 0)); - HIP_CHECK(hipHostRegister(&(dataPtr.get()[1]), sizeof(intT), 0)); + // Cleanup + HIP_CHECK(hipStreamDestroy(stream)); +} - // Allocate Signal Memory - int64_t* signalPtr; - HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory)); +template +void syncAndCheckData(hipStream_t stream, UIntT* dataPtr, TestPtr signalPtr, size_t offset, + TEST_WAIT tc, std::array& events) { + // Ensure first part of host memory is updated + HIP_CHECK(hipEventSynchronize(events[0])); + REQUIRE(dataPtr[0] == DATA_UPDATE); - std::vector events; - initData(dataPtr.get(), signalPtr, tc, events); - - if (std::is_same::value) { - CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag) - HIP_CHECK(hipEventRecord(events[0], stream)); - - if (static_cast(tc.mask) != 0xFFFFFFFF) { - CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, static_cast(tc.waitValue), - tc.compareOp, static_cast(tc.mask)); - } else { - CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, tc.waitValue, tc.compareOp); - } - - CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag) - } else { - CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag) - HIP_CHECK(hipEventRecord(events[0], stream)); - - if (tc.mask != 0xFFFFFFFFFFFFFFFF) { - CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp, tc.mask); - } else { - CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp); - } - - CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag) + if constexpr (isBlocking) { + // Ensure second part of host memory isn't updated yet + HIP_CHECK_ERROR(hipEventQuery(events[1]), hipErrorNotReady); + REQUIRE(dataPtr[1] == DATA_INIT); + // Update value to release stream + signalPtr.setValue(tc.signalValuePass, offset); } + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipEventQuery(events[1])); + // Finally ensure that second part of host memory is updated + REQUIRE(dataPtr[1] == DATA_UPDATE); +} + + +template +void testWait(TEST_WAIT tc) { + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-128"); + return; +#endif + + using UIntT = typename TestType::UIntType; + constexpr auto ptrType = TestType::ptrType; + constexpr UIntT defaultMask = ~static_cast(0); + + // Initialize stream + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocate Host Memory + auto dataPtr = std::make_unique(2); + // Register Host Memory + HIP_CHECK(hipHostRegister(dataPtr.get(), sizeof(UIntT), 0)); + HIP_CHECK(hipHostRegister(dataPtr.get() + 1, sizeof(UIntT), 0)); + std::fill(dataPtr.get(), dataPtr.get() + 2, DATA_INIT); + + std::array events; + HIP_CHECK(hipEventCreate(&events[0])); + HIP_CHECK(hipEventCreate(&events[1])); + + + const auto offsets = get_offsets(); + const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets))); + + auto waitPtr = allocMem(); + UIntT* const target = waitPtr.ptr + offset; + waitPtr.setValue(isBlocking ? tc.signalValueFail : tc.signalValuePass, offset); + + HIP_CHECK(writeFunc(stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)); + HIP_CHECK(hipEventRecord(events[0], stream)); + + if (tc.mask != defaultMask) { + HIP_CHECK(waitFunc(stream, target, tc.waitValue, tc.compareOp, tc.mask)); + } else { + HIP_CHECK(waitFunc(stream, target, tc.waitValue, tc.compareOp)); + } + + HIP_CHECK(writeFunc(stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)); + HIP_CHECK(hipEventRecord(events[1], stream)); - syncAndCheckData(stream, dataPtr.get(), signalPtr, tc, events); - cleanup(stream, dataPtr.get(), signalPtr); -} -#undef CHECK_HIP_STREAM_VALUE + syncAndCheckData(stream, dataPtr.get(), std::move(waitPtr), offset, tc, events); + // Cleanup + HIP_CHECK(hipEventDestroy(events[0])); + HIP_CHECK(hipEventDestroy(events[1])); + HIP_CHECK(hipHostUnregister(dataPtr.get())); + HIP_CHECK(hipHostUnregister(dataPtr.get() + 1)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead #define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32(suffix, test_t) \ TEST_CASE("Unit_hipStreamValue_Wait32_Blocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ } \ TEST_CASE("Unit_hipStreamValue_Wait32_NonBlocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ } + // Using Mask -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_1", - TEST_WAIT64( // mask will ignore few MSB bits - hipStreamWaitValueGte, 0x0000FFFFFFFFFFFF, - 0x000000007FFF0001, 0x7FFF00007FFF0000, - 0x000000007FFF0001)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_2", - TEST_WAIT64(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6)) +DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte", + TEST_WAIT32(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_1", - TEST_WAIT64( // mask will ignore few MSB bits - hipStreamWaitValueEq, 0x0000FFFFFFFFFFFF, - 0x000000000FFF0001, 0x7FFF00000FFF0000, - 0x7F0000000FFF0001)) + TEST_WAIT32( // mask will ignore few MSB bits + hipStreamWaitValueEq, 0x0000FFFF, 0x00000001, + 0x0FFF0000, 0x0FFF0001)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_2", - TEST_WAIT64(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11)) + TEST_WAIT32(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_And", - TEST_WAIT64( // mask will discard bits 8 to 11 + TEST_WAIT32( // mask will discard bits 8 to 11 hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_1", - TEST_WAIT64( // mask is set to ignore the sign bit. - hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF, - 0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd, - 0x7FFFFFFFFFFFFdb5)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_2", - TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3. - hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76)) // Not Using Mask DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Eq", @@ -299,19 +400,47 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_And", TEST_WAIT32(hipStreamWaitValueAnd, 0x70F0F0F0, 0x0F0F0F0F, 0X1F0F0F0F)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Nor", - TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, - static_cast(0x85555555), - static_cast(0x9AAAAAAA))) + TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, 0x85555555, + 0x9AAAAAAA)) #undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32 +#if HT_AMD +// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead #define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \ TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ + SECTION("Signal") { testWait, true>(test_t); } \ } \ TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ + SECTION("Signal") { testWait, false>(test_t); } \ } +#else +#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \ + TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ + } \ + TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ + } +#endif // Using Mask @@ -332,14 +461,6 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Eq_2", DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_And", TEST_WAIT64( // mask will discard bits 8 to 11 hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_1", - TEST_WAIT64( // mask is set to ignore the sign bit. - hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF, - 0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd, - 0x7FFFFFFFFFFFFdb5)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_2", - TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3. - hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Gte", TEST_WAIT64(hipStreamWaitValueGte, 0x7FFFFFFFFFFF0001, @@ -352,94 +473,72 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_And", 0x0F0F0F0F0F0F0F0F, 0X1F0F0F0F0F0F0F0F)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Nor", TEST_WAIT64(hipStreamWaitValueNor, 0x4724724747247247, - static_cast(0xbddbddbdbddbddbd), - static_cast(0xbddbddbdbddbddb3))) + 0xbddbddbdbddbddbd, 0xbddbddbdbddbddb3)) #undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64 -#endif - // Negative Tests TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") { - #if HT_AMD HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96"); return; #endif + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); - REQUIRE(stream != nullptr); - // Allocate Host Memory - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); - - // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); - - // Set dummy data - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; - - auto compareOp = hipStreamWaitValueGte; + const auto compareOp = hipStreamWaitValueGte; + const auto expectedError = hipErrorInvalidValue; // Memory pointer negative tests - - INFO("Testing Invalid Memory Pointer for hipStreamWriteValue32"); - NEG_TEST_ERROR_CHECK(Write, 32, hipErrorNotSupported, stream, nullptr, 0, writeFlag) - - INFO("Testing Invalid Memory Pointer for hipStreamWriteValue64"); - NEG_TEST_ERROR_CHECK(Write, 64, hipErrorNotSupported, stream, nullptr, 0, writeFlag) - - INFO("Testing Invalid Memory Pointer for hipStreamWaitValue32"); - NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, nullptr, 0, compareOp) - - INFO("Testing Invalid Memory Pointer for hipStreamWaitValue64"); - NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, nullptr, 0, compareOp) + SECTION("Invalid Memory Pointer for hipStreamWriteValue32") { + HIP_CHECK_ERROR(hipStreamWriteValue32(stream, nullptr, 0, writeFlag), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWriteValue64") { + HIP_CHECK_ERROR(hipStreamWriteValue64(stream, nullptr, 0, writeFlag), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWaitValue32") { + HIP_CHECK_ERROR(hipStreamWaitValue32(stream, nullptr, 0, compareOp), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWaitValue32") { + HIP_CHECK_ERROR(hipStreamWaitValue64(stream, nullptr, 0, compareOp), expectedError); + } // Cleanup - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); HIP_CHECK(hipStreamDestroy(stream)); } -TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag") { +TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t, uint64_t) { #if HT_AMD HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96"); return; #endif + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); - REQUIRE(stream != nullptr); // Allocate Host Memory - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); + auto hostPtr = std::make_unique(); // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); + HIP_CHECK(hipHostRegister(hostPtr.get(), sizeof(TestType), 0)); // Set dummy data - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; + *hostPtr = 0x0; /* EXSWCPHIPT-96 */ - INFO("Testing Invalid flag for hipStreamWaitValue32"); - NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, hostPtr32.get(), 0, -1) - INFO("Testing Invalid flag for hipStreamWaitValue64"); - NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, hostPtr64.get(), 0, -1) + HIP_CHECK_ERROR(waitFunc(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue); // Cleanup - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); + HIP_CHECK(hipHostUnregister(hostPtr.get())); HIP_CHECK(hipStreamDestroy(stream)); } - -#undef NEG_TEST_ERROR_CHECK diff --git a/catch/unit/stream/hipStreamWithCUMask.cc b/catch/unit/stream/hipStreamWithCUMask.cc index b65cdc4ca1..ea6b033d9c 100644 --- a/catch/unit/stream/hipStreamWithCUMask.cc +++ b/catch/unit/stream/hipStreamWithCUMask.cc @@ -149,8 +149,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") { HIP_CHECK(hipGetDeviceProperties(&props, 0)); createDefaultCUMask(&defaultCUMask, props.multiProcessorCount); - hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(), - defaultCUMask.data()); + HIP_CHECK(hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(), + defaultCUMask.data())); HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); const unsigned blocks = GRIDSIZE; @@ -244,7 +244,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, streams[0], dA[0], dC[0], N); - hipDeviceSynchronize(); + HIP_CHECK(hipDeviceSynchronize()); auto single_end = std::chrono::steady_clock::now(); std::chrono::duration single_kernel_time = single_end - single_start; @@ -264,7 +264,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, streams[np], dA[np], dC[np], N); } - hipDeviceSynchronize(); + HIP_CHECK(hipDeviceSynchronize()); auto all_end = std::chrono::steady_clock::now(); std::chrono::duration all_kernel_time = all_end - all_start; @@ -288,8 +288,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { delete [] hA; delete [] hC; for (int np = 0; np < KNumPartition; np++) { - hipFree(dC[np]); - hipFree(dA[np]); + HIP_CHECK(hipFree(dC[np])); + HIP_CHECK(hipFree(dA[np])); HIP_CHECK(hipStreamDestroy(streams[np])); } }