diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index ceaabe158f..67e11ec779 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -41,10 +41,6 @@ "Unit_hipStreamCreateWithFlags_DefaultStreamInteraction", "Unit_hipStreamWaitEvent_UninitializedStream_Negative", "Unit_hipDeviceSetSharedMemConfig_Negative_Parameters", - "Disabling tests tracked with SWDEV-394083", - "Unit_hipDeviceSynchronize_Positive_Nullstream", - "Disabling tests tracked with SWDEV-393637", - "Unit_hipDeviceSynchronize_Functional", "Unit_hipMemset3DSync", "Unit_hipStreamAddCallback_StrmSyncTiming", "Disabling test tracked SWDEV-394199", diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 2ba1dddba3..7f9e859cd0 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -22,6 +22,7 @@ THE SOFTWARE. #pragma once #include "hip_test_context.hh" + #include #include #include @@ -30,6 +31,7 @@ THE SOFTWARE. #include #include #include +#include #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); @@ -429,6 +431,52 @@ static inline void runKernelForDuration(std::chrono::milliseconds duration, hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000); } +class BlockingContext { + std::atomic_bool blocked{true}; + hipStream_t stream; + + public: + BlockingContext(hipStream_t s) : stream(s), blocked(true) {} + + BlockingContext(const BlockingContext& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + } + + BlockingContext(const BlockingContext&& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + } + + void reset() { blocked = true; } + + BlockingContext& operator=(const BlockingContext& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + return *this; + } + + void block_stream() { + blocked = true; + auto blocking_callback = [](hipStream_t, hipError_t, void* data) { + auto blocked = reinterpret_cast(data); + while (blocked->load()) { + // Yield this thread till we are waiting + std::this_thread::yield(); + } + }; + HIP_CHECK(hipStreamAddCallback(stream, blocking_callback, (void*)&blocked, 0)); + } + + void unblock_stream() { + blocked = false; + } + + bool is_blocked() const { return hipStreamQuery(stream) == hipErrorNotReady; } + + bool blocked_val() const { return blocked.load(); } + hipStream_t stream_val() const { return stream; } +}; } // namespace HipTest // This must be called in the beginning of image test app's main() to indicate whether image diff --git a/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc b/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc index bfb02ab784..4dd3ed0dff 100644 --- a/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc +++ b/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc @@ -33,15 +33,15 @@ THE SOFTWARE. #define NUM_ITERS 1 << 30 static __global__ void Iter(int* Ad, int num) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - // Kernel loop designed to execute very slowly. - // so we can test timing-related - // behavior below - if (tx == 0) { - for (int i = 0; i < num; i++) { - Ad[tx] += 1; - } + int tx = threadIdx.x + blockIdx.x * blockDim.x; + // Kernel loop designed to execute very slowly. + // so we can test timing-related + // behavior below + if (tx == 0) { + for (int i = 0; i < num; i++) { + Ad[tx] += 1; } + } } TEST_CASE("Unit_hipDeviceSynchronize_Positive_Empty_Streams") { @@ -61,42 +61,47 @@ TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") { INFO("Current device: " << device); int *A_h = nullptr, *A_d = nullptr; + HipTest::BlockingContext b_context{nullptr}; HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), _SIZE, hipHostMallocDefault)); A_h[0] = 1; HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), _SIZE)); HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL)); + b_context.block_stream(); + REQUIRE(b_context.is_blocked()); hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30); HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL)); - CHECK(1 << 30 != A_h[0] - 1); + REQUIRE(1 << 30 != A_h[0] - 1); + b_context.unblock_stream(); HIP_CHECK(hipDeviceSynchronize()); - CHECK(1 << 30 == A_h[0] - 1); + REQUIRE(1 << 30 == A_h[0] - 1); } TEST_CASE("Unit_hipDeviceSynchronize_Functional") { int* A[NUM_STREAMS]; int* Ad[NUM_STREAMS]; hipStream_t stream[NUM_STREAMS]; + std::vector b_context; + b_context.reserve(NUM_STREAMS); + for (int i = 0; i < NUM_STREAMS; i++) { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A[i]), _SIZE, - hipHostMallocDefault)); - A[i][0] = 1; - HIP_CHECK(hipMalloc(reinterpret_cast(&Ad[i]), _SIZE)); - HIP_CHECK(hipStreamCreate(&stream[i])); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A[i]), _SIZE, hipHostMallocDefault)); + A[i][0] = 1; + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad[i]), _SIZE)); + HIP_CHECK(hipStreamCreate(&stream[i])); + b_context.emplace_back(HipTest::BlockingContext(stream[i])); } for (int i = 0; i < NUM_STREAMS; i++) { - HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, - stream[i])); + HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i])); } for (int i = 0; i < NUM_STREAMS; i++) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, - stream[i], Ad[i], NUM_ITERS); - HIP_CHECK(hipGetLastError()); + b_context[i].block_stream(); + REQUIRE(b_context[i].is_blocked()); + hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], NUM_ITERS); } for (int i = 0; i < NUM_STREAMS; i++) { - HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, - stream[i])); + HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i])); } @@ -106,7 +111,10 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") { // Conservative implementations which synchronize the hipMemcpyAsync will // fail, ie if HIP_LAUNCH_BLOCKING=true. - CHECK(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1); + REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1); + for (int i = 0; i < NUM_STREAMS; i++) { + b_context[i].unblock_stream(); + } HIP_CHECK(hipDeviceSynchronize()); - CHECK(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1); + REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1); }