From 8c62d0fbc42ec0cf2cabc24187ec8e936c554ab9 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 10 Nov 2017 00:14:34 +0000 Subject: [PATCH] Update new tests so as to make them work with new variadic based launch mechanisms. --- .../src/runtimeApi/memory/hipMemcpyDtoD.cpp | 34 ++++++++++++++----- .../runtimeApi/memory/hipMemcpyDtoDAsync.cpp | 34 ++++++++++++++----- .../src/runtimeApi/memory/hipMemcpyPeer.cpp | 30 ++++++++++++---- .../runtimeApi/memory/hipMemcpyPeerAsync.cpp | 32 +++++++++++++---- 4 files changed, 101 insertions(+), 29 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp index ccb02b74ce..c64b01f8a7 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp @@ -49,21 +49,39 @@ int main() HIPCHECK(hipMalloc(&Y_d,Nbytes)); HIPCHECK(hipMalloc(&Z_d,Nbytes)); - + HIPCHECK(hipSetDevice(0)); HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(A_d), + static_cast(B_d), + C_d, + N); HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - - + + HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes)); + HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes)); HIPCHECK(hipMemcpyDtoD(Y_d, B_d, Nbytes)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(X_d), + static_cast(Y_d), + Z_d, + N); HIPCHECK(hipMemcpyDtoH(C_h, Z_d, Nbytes)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); @@ -73,8 +91,8 @@ int main() HIPCHECK(hipFree(Y_d)); HIPCHECK(hipFree(Z_d)); } - + passed(); - + } diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp index 5c99b43564..6d21ac62e7 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp @@ -50,25 +50,43 @@ int main() HIPCHECK(hipMalloc(&Y_d,Nbytes)); HIPCHECK(hipMalloc(&Z_d,Nbytes)); - + HIPCHECK(hipSetDevice(0)); HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(A_d), + static_cast(B_d), + C_d, + N); HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - - HIPCHECK(hipStreamCreate(&s)); + + HIPCHECK(hipStreamCreate(&s)); HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s)); + HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s)); HIPCHECK(hipMemcpyDtoDAsync(Y_d, B_d, Nbytes, s)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(X_d), + static_cast(Y_d), + Z_d, + N); HIPCHECK(hipMemcpyDtoHAsync(C_h, Z_d, Nbytes, s)); HIPCHECK(hipStreamSynchronize(s)); HIPCHECK(hipDeviceSynchronize()); - + HipTest::checkVectorADD(A_h, B_h, C_h, N); HIPCHECK(hipStreamDestroy(s)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); @@ -78,7 +96,7 @@ int main() } passed(); - + } diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp index 7e2fc2d3d0..95b19c1090 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp @@ -48,24 +48,42 @@ int main() HIPCHECK(hipMalloc(&X_d,Nbytes)); HIPCHECK(hipMalloc(&Y_d,Nbytes)); HIPCHECK(hipMalloc(&Z_d,Nbytes)); - + HIPCHECK(hipSetDevice(0)); HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(A_d), + static_cast(B_d), + C_d, + N); HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - + HIPCHECK(hipSetDevice(1)); hipMemcpyPeer(X_d, 1, A_d, 0, Nbytes); //this call is eqv to hipMemcpy(hipMemcpyD2D) which goes via stg bufs. hipMemcpyPeer(Y_d, 1, B_d, 0, Nbytes); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(X_d), + static_cast(Y_d), + Z_d, + N); HIPCHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIPCHECK(hipFree(X_d)); HIPCHECK(hipFree(Y_d)); @@ -74,7 +92,7 @@ int main() passed(); - + } diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp index 9d46ccb0d8..943e4a6b95 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp @@ -51,26 +51,44 @@ int main() HIPCHECK(hipMalloc(&Y_d,Nbytes)); HIPCHECK(hipMalloc(&Z_d,Nbytes)); - + HIPCHECK(hipSetDevice(0)); HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(A_d), + static_cast(B_d), + C_d, + N); HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK (hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - - HIPCHECK(hipStreamCreate(&s)); + + HIPCHECK(hipStreamCreate(&s)); HIPCHECK(hipSetDevice(1)); HIPCHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, Nbytes, s)); HIPCHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, Nbytes, s)); - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + hipLaunchKernel( + HipTest::vectorADD, + dim3(blocks), + dim3(threadsPerBlock), + 0, + 0, + static_cast(X_d), + static_cast(Y_d), + Z_d, + N); HIPCHECK ( hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost)); HIPCHECK (hipDeviceSynchronize()); HIPCHECK (hipStreamSynchronize(s)); HipTest::checkVectorADD(A_h, B_h, C_h, N); - + HIPCHECK(hipStreamDestroy(s)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIPCHECK(hipFree(X_d)); @@ -79,7 +97,7 @@ int main() } passed(); - + }