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

SWDEV-314080 - Tested All hipMemcpy**() apis with hipStreamPerThread stream obj (#2496)

Browse Source 
Change-Id: I8f429eb0cc3be2e4d62c76ccb8c1510c56a1e143
This commit is contained in:
ROCm CI Service Account
2022-02-24 18:58:04 +05:30
committed by GitHub
parent 7d8d41e45b
commit d18b4b8848
7 changed files with 219 additions and 29 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/catch/unit/memory/hipMemcpy2DAsync.cc
+31 -13
View File
@@ -83,22 +83,40 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
// Initialize the data
HipTest::setDefaultData<TestType>(NUM_W*NUM_H, A_h, B_h, C_h);
SECTION("Calling Async apis with stream object created by user") {
// Host to Device
HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyHostToDevice, stream));
// Host to Device
HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyHostToDevice, stream));
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice, stream));
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice, stream));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling Async apis with hipStreamPerThread") {
// Host to Device
HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyHostToDevice, hipStreamPerThread));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice, hipStreamPerThread));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyDeviceToHost, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
// Validating the result
REQUIRE(HipTest::checkArray<TestType>(A_h, B_h, COLUMNS, ROWS) == true);
tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc
+12 -4
View File
@@ -62,10 +62,18 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") {
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") {
HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") {
HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true);
// Cleaning the memory
tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc
+12 -5
View File
@@ -58,11 +58,18 @@ TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") {
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") {
HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") {
HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost));
tests/catch/unit/memory/hipMemcpy3DAsync.cc
+8 -2
View File
@@ -602,8 +602,14 @@ void Memcpy3DAsync<T>::simple_Memcpy3DAsync() {
#else
myparms.kind = hipMemcpyHostToDevice;
#endif
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
SECTION("Calling hipMemcpy3DAsync() using user declared stream obj") {
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpy3DAsync() using hipStreamPerThread") {
REQUIRE(hipMemcpy3DAsync(&myparms, hipStreamPerThread) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
// Array to Array
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
tests/catch/unit/memory/hipMemcpyPeerAsync.cc
+14 -5
View File
@@ -149,11 +149,20 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
// Copying data from GPU-0 to GPU-1 and performing vector addition
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
stream));
HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
stream));
HIP_CHECK(hipStreamSynchronize(stream));
SECTION("Calling hipMemcpyPerAsync() using user defined stream obj") {
HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
stream));
HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") {
HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
hipStreamPerThread));
HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(X_d),
static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));