SWDEV-393637 - Standardize the way to implement wait mechanism (#316)

Change-Id: I59027667806878191f9c641f27ce47e6f85ba40c
This commit is contained in:
ROCm CI Service Account authored and GitHub committed 2023-06-21 21:27:36 +05:30
1 والد efc742a900
کامیت 5f4af7cde4
3 files changed
+80 -28

No files matched your search

@@ -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<void**>(&A_h), _SIZE, hipHostMallocDefault));
A_h[0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&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<HipTest::BlockingContext> b_context;
b_context.reserve(NUM_STREAMS);
for (int i = 0; i < NUM_STREAMS; i++) {
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE,
hipHostMallocDefault));
A[i][0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
HIP_CHECK(hipStreamCreate(&stream[i]));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE, hipHostMallocDefault));
A[i][0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&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);
}