diff --git a/hipamd/tests/src/hipStreamL5.cpp b/hipamd/tests/src/hipStreamL5.cpp index 5a65dc50dc..b8dfa120d0 100644 --- a/hipamd/tests/src/hipStreamL5.cpp +++ b/hipamd/tests/src/hipStreamL5.cpp @@ -36,6 +36,7 @@ For example, a test for Ker, D2D, D2H, H2H, H2D is given as test34512(); Note that all memory copies are Async. +Invalid { *WARNING: The commented out assertions are failing cases. According to my observation, they are happening with tests which end in HostToHost and take data from previous @@ -49,6 +50,7 @@ For disjoint data passed: test24513 test25134 test34512 +} */ template @@ -199,633 +201,632 @@ void test15234(){ template void test23451(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = sizeof(T) * N; +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); +const size_t N = 1000; +const size_t size = sizeof(T) * N; - T *Ah, *Bh, *Ch; - T *Ad, *Bd; +T *Ah, *Bh, *Ch; +T *Ad, *Bd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); - setArray(Ah, N, T(1)); +setArray(Ah, N, T(1)); - H2DAsync(Ad, Ah, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Bh, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); - HIPCHECK(hipDeviceSynchronize()); - //HIPASSERT(Ah[10] == Ch[10]); +H2DAsync(Ad, Ah, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Bh, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); +HIPCHECK(hipDeviceSynchronize()); +HIPASSERT(Ah[10] + T(1) == Ch[10]); } template void test24513(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = sizeof(T) * N; +const size_t N = 1000; +const size_t size = sizeof(T) * N; - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Ah, N, T(1)); - setArray(Dh, N, T(2)); +setArray(Ah, N, T(1)); +setArray(Dh, N, T(2)); - H2D(Cd, Dh, size); +H2D(Cd, Dh, size); - H2DAsync(Ad, Ah, size, stream); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Bh, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); - HIPCHECK(hipDeviceSynchronize()); +H2DAsync(Ad, Ah, size, stream); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Bh, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); +HIPCHECK(hipDeviceSynchronize()); - D2H(Eh, Cd, size); +D2H(Eh, Cd, size); - HIPASSERT(Eh[0] == Dh[0] + T(1)); - //HIPASSERT(Ah[0] == Ch[0]); +HIPASSERT(Eh[0] == Dh[0] + T(1)); +HIPASSERT(Ah[0] == Ch[0]); } template void test25134(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = sizeof(T) * N; +const size_t N = 1000; +const size_t size = sizeof(T) * N; - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Ah, N, T(1)); - setArray(Dh, N, T(2)); +setArray(Ah, N, T(1)); +setArray(Dh, N, T(2)); - H2D(Bd, Dh, size); +H2D(Bd, Dh, size); - H2DAsync(Ad, Ah, size, stream); - D2HAsync(Bh, Ad, size, stream); - H2HAsync(Ch, Bh, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); - D2DAsync(Cd, Bd, size, stream); +H2DAsync(Ad, Ah, size, stream); +D2HAsync(Bh, Ad, size, stream); +H2HAsync(Ch, Bh, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); +D2DAsync(Cd, Bd, size, stream); - D2H(Eh, Cd, size); +D2H(Eh, Cd, size); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); + +HIPASSERT(Ah[10] == Ch[10]); +HIPASSERT(Dh[10] + T(1) == Eh[10]); - //HIPASSERT(Ah[10] == Ch[10]); - HIPASSERT(Dh[10] + T(1) == Eh[10]); } template void test21345(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch, *Dh; - T *Ad, *Bd; +T *Ah, *Bh, *Ch, *Dh; +T *Ad, *Bd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, true); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, true); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); - setArray(Ah, N, T(1)); - setArray(Bh, N, T(2)); +setArray(Ah, N, T(1)); +setArray(Bh, N, T(2)); - H2DAsync(Ad, Ah, size, stream); - H2HAsync(Ch, Bh, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Dh, Bd, size, stream); +H2DAsync(Ad, Ah, size, stream); +H2HAsync(Ch, Bh, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Dh, Bd, size, stream); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT( Bh[10] == Ch[10] ); - HIPASSERT( Ah[10] + T(1) == Dh[10]); +HIPASSERT( Bh[10] == Ch[10] ); +HIPASSERT( Ah[10] + T(1) == Dh[10]); } template void test34512(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Bh, *Ch, *Dh; - T *Ah, *Eh; - T *Ad, *Bd, *Cd; +T *Bh, *Ch, *Dh; +T *Ah, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, true); - initArrays(&Ah, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, true); +initArrays(&Ah, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Ah, N, T(1)); +setArray(Ah, N, T(1)); - H2D(Ad, Ah, size); +H2D(Ad, Ah, size); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Bh, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); - H2DAsync(Cd, Ch, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Bh, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); +H2DAsync(Cd, Ch, size, stream); - D2H(Dh, Cd, size); +D2H(Dh, Cd, size); - HIPCHECK(hipDeviceSynchronize()); - //HIPASSERT( Ah[10] + T(1) == Dh[10] ); +HIPCHECK(hipDeviceSynchronize()); +HIPASSERT( Ah[10] + T(1) == Dh[10] ); } template void test35124(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh; - T *Ch, *Dh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh; +T *Ch, *Dh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, false); - initArrays(&Dh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, false); +initArrays(&Dh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(1)); +setArray(Dh, N, T(1)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - D2HAsync(Ah, Ad, size, stream); - H2HAsync(Bh, Ah, size, stream); - H2DAsync(Bd, Bh, size, stream); - D2DAsync(Cd, Bd, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +D2HAsync(Ah, Ad, size, stream); +H2HAsync(Bh, Ah, size, stream); +H2DAsync(Bd, Bh, size, stream); +D2DAsync(Cd, Bd, size, stream); - D2H(Ch, Cd, size); +D2H(Ch, Cd, size); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Dh[10] + T(1) == Ch[10]); +HIPASSERT(Dh[10] + T(1) == Ch[10]); } template void test31245(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +const size_t N = 1000; +const size_t size = N * sizeof(T); +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(1)); - setArray(Ah, N, T(2)); +setArray(Dh, N, T(1)); +setArray(Ah, N, T(2)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - H2HAsync(Bh, Ah, size, stream); - H2DAsync(Bd, Bh, size, stream); - D2DAsync(Cd, Bd, size, stream); - D2HAsync(Ch, Cd, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +H2HAsync(Bh, Ah, size, stream); +H2DAsync(Bd, Bh, size, stream); +D2DAsync(Cd, Bd, size, stream); +D2HAsync(Ch, Cd, size, stream); - D2H(Eh, Ad, size); +D2H(Eh, Ad, size); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Dh[10] + T(1) == Eh[10]); - HIPASSERT(Bh[10] == Ch[10]); +HIPASSERT(Dh[10] + T(1) == Eh[10]); +HIPASSERT(Bh[10] == Ch[10]); } template void test32451(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Ah, N, T(1)); - setArray(Eh, N, T(2)); +setArray(Ah, N, T(1)); +setArray(Eh, N, T(2)); - H2D(Ad, Eh, size); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); - H2DAsync(Bd, Ah, size, stream); - D2DAsync(Cd, Bd, size, stream); - D2HAsync(Bh, Cd, size, stream); - H2HAsync(Ch, Bh, size, stream); - HIPCHECK(hipDeviceSynchronize()); - D2H(Dh, Ad, size); +H2D(Ad, Eh, size); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad); +H2DAsync(Bd, Ah, size, stream); +D2DAsync(Cd, Bd, size, stream); +D2HAsync(Bh, Cd, size, stream); +H2HAsync(Ch, Bh, size, stream); +D2H(Dh, Ad, size); +HIPCHECK(hipDeviceSynchronize()); - //HIPASSERT(Ah[10] == Ch[10]); - HIPASSERT(Eh[10] + T(1) == Dh[10]); +HIPASSERT(Ah[10] == Ch[10]); +HIPASSERT(Eh[10] + T(1) == Dh[10]); } template void test45123(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh; - T *Ch, *Dh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh; +T *Ch, *Dh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, false); - initArrays(&Dh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, false); +initArrays(&Dh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(1)); +setArray(Dh, N, T(1)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Ah, Bd, size, stream); - H2HAsync(Bh, Ah, size, stream); - H2DAsync(Cd, Bh, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); - D2H(Ch, Cd, size); - HIPCHECK(hipDeviceSynchronize()); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Ah, Bd, size, stream); +H2HAsync(Bh, Ah, size, stream); +H2DAsync(Cd, Bh, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); +D2H(Ch, Cd, size); +HIPCHECK(hipDeviceSynchronize()); + +HIPASSERT(Dh[10] + T(1) == Ch[10]); - HIPASSERT(Dh[10] + T(1) == Ch[10]); } template void test41235(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh; - T *Ch; - T *Ad, *Bd, *Cd; +T *Ah, *Bh; +T *Ch; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Ch, N, T(1)); +setArray(Ch, N, T(1)); - H2D(Ad, Ch, size); +H2D(Ad, Ch, size); - D2DAsync(Bd, Ad, size, stream); - D2HAsync(Ah, Bd, size, stream); - H2DAsync(Cd, Ah, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); - D2HAsync(Bh, Cd, size, stream); +D2DAsync(Bd, Ad, size, stream); +D2HAsync(Ah, Bd, size, stream); +H2DAsync(Cd, Ah, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); +D2HAsync(Bh, Cd, size, stream); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Ch[10] + T(1) == Bh[10]); +HIPASSERT(Ch[10] + T(1) == Bh[10]); } template void test42351(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(2)); - setArray(Ah, N, T(1)); +setArray(Dh, N, T(2)); +setArray(Ah, N, T(1)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - D2DAsync(Bd, Ad, size, stream); - H2DAsync(Cd, Ah, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); - D2HAsync(Bh, Cd, size, stream); - H2HAsync(Ch, Bh, size, stream); +D2DAsync(Bd, Ad, size, stream); +H2DAsync(Cd, Ah, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd); +D2HAsync(Bh, Cd, size, stream); +H2HAsync(Ch, Bh, size, stream); - D2H(Eh, Bd, size); +D2H(Eh, Bd, size); + +HIPCHECK(hipDeviceSynchronize()); +HIPASSERT(Dh[10] == Eh[10]); +HIPASSERT(Ah[10] + T(1) == Ch[10]); - HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Dh[10] == Eh[10]); - //HIPASSERT(Ah[10] + T(1) == Ch[10]); } template void test43512(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh; - T *Ch, *Dh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh; +T *Ch, *Dh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, false); - initArrays(&Dh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, false); +initArrays(&Dh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(1)); +setArray(Dh, N, T(1)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - D2DAsync(Bd, Ad, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); - D2HAsync(Ah, Bd, size, stream); - H2HAsync(Bh, Ah, size, stream); - H2DAsync(Cd, Bh, size, stream); +D2DAsync(Bd, Ad, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); +D2HAsync(Ah, Bd, size, stream); +H2HAsync(Bh, Ah, size, stream); +H2DAsync(Cd, Bh, size, stream); - D2H(Ch, Cd, size); - HIPCHECK(hipDeviceSynchronize()); - //HIPASSERT( Dh[10] + T(1) == Ch[10]); +D2H(Ch, Cd, size); +HIPCHECK(hipDeviceSynchronize()); +HIPASSERT( Dh[10] + T(1) == Ch[10]); } template void test51234(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh; - T *Ch, *Dh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh; +T *Ch, *Dh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, false); - initArrays(&Dh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, false); +initArrays(&Dh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Dh, N, T(1)); +setArray(Dh, N, T(1)); - H2D(Ad, Dh, size); +H2D(Ad, Dh, size); - D2HAsync(Ah, Ad, size, stream); - H2HAsync(Bh, Ah, size, stream); - H2DAsync(Bd, Bh, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); - D2DAsync(Cd, Bd, size, stream); +D2HAsync(Ah, Ad, size, stream); +H2HAsync(Bh, Ah, size, stream); +H2DAsync(Bd, Bh, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); +D2DAsync(Cd, Bd, size, stream); - D2H(Ch, Cd, size); +D2H(Ch, Cd, size); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Ch[10] == Dh[10] + T(1)); +HIPASSERT(Ch[10] == Dh[10] + T(1)); } template void test52341(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch; - T *Dh, *Eh; - T *Ad, *Bd, *Cd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh; +T *Ad, *Bd, *Cd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); - setArray(Eh, N, T(1)); - setArray(Bh, N, T(2)); +setArray(Eh, N, T(1)); +setArray(Bh, N, T(2)); - H2D(Ad, Eh, size); +H2D(Ad, Eh, size); - D2HAsync(Ah, Ad, size, stream); - H2DAsync(Bd, Ah, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); - D2DAsync(Cd, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); +D2HAsync(Ah, Ad, size, stream); +H2DAsync(Bd, Ah, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); +D2DAsync(Cd, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); - D2H(Dh, Cd, size); +D2H(Dh, Cd, size); - HIPCHECK(hipDeviceSynchronize()); +HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Eh[10] + T(1) == Dh[10]); - HIPASSERT(Ch[10] == Bh[10]); +HIPASSERT(Eh[10] + T(1) == Dh[10]); +HIPASSERT(Ch[10] == Bh[10]); } template void test53412(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = sizeof(T) * N; +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); +const size_t N = 1000; +const size_t size = sizeof(T) * N; - T *Ah, *Bh, *Ch, *Dh; - T *Eh, *Fh, *Gh; - T *Ad, *Bd, *Cd, *Dd; +T *Ah, *Bh, *Ch, *Dh; +T *Eh, *Fh, *Gh; +T *Ad, *Bd, *Cd, *Dd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, true); - initArrays(&Eh, N, false, false); - initArrays(&Fh, N, false, false); - initArrays(&Gh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); - initArrays(&Dd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, true); +initArrays(&Eh, N, false, false); +initArrays(&Fh, N, false, false); +initArrays(&Gh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); +initArrays(&Dd, N, true, false); - setArray(Dh, N, T(1)); - setArray(Eh, N, T(2)); - setArray(Bh, N, T(3)); +setArray(Dh, N, T(1)); +setArray(Eh, N, T(2)); +setArray(Bh, N, T(3)); - H2D(Ad, Dh, size); - H2D(Bd, Eh, size); +H2D(Ad, Dh, size); +H2D(Bd, Eh, size); - D2HAsync(Ah, Ad, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); - D2DAsync(Cd, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); - H2DAsync(Dd, Ch, size, stream); +D2HAsync(Ah, Ad, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd); +D2DAsync(Cd, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); +H2DAsync(Dd, Ch, size, stream); - D2H(Fh, Cd, size); - D2H(Gh, Dd, size); +D2H(Fh, Cd, size); +D2H(Gh, Dd, size); - HIPASSERT(Ah[10] == Dh[10]); - HIPASSERT(Eh[10] + T(1) == Fh[10]); - HIPASSERT(Bh[10] == Gh[10]); +HIPASSERT(Ah[10] == Dh[10]); +HIPASSERT(Eh[10] + T(1) == Fh[10]); +HIPASSERT(Bh[10] == Gh[10]); } template void test54123(){ - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); +hipStream_t stream; +HIPCHECK(hipStreamCreate(&stream)); - const size_t N = 1000; - const size_t size = N * sizeof(T); +const size_t N = 1000; +const size_t size = N * sizeof(T); - T *Ah, *Bh, *Ch; - T *Dh, *Eh, *Fh, *Gh; - T *Ad, *Bd, *Cd, *Dd; +T *Ah, *Bh, *Ch; +T *Dh, *Eh, *Fh, *Gh; +T *Ad, *Bd, *Cd, *Dd; - initArrays(&Ah, N, false, true); - initArrays(&Bh, N, false, true); - initArrays(&Ch, N, false, true); - initArrays(&Dh, N, false, false); - initArrays(&Eh, N, false, false); - initArrays(&Fh, N, false, false); - initArrays(&Gh, N, false, false); - initArrays(&Ad, N, true, false); - initArrays(&Bd, N, true, false); - initArrays(&Cd, N, true, false); - initArrays(&Dd, N, true, false); +initArrays(&Ah, N, false, true); +initArrays(&Bh, N, false, true); +initArrays(&Ch, N, false, true); +initArrays(&Dh, N, false, false); +initArrays(&Eh, N, false, false); +initArrays(&Fh, N, false, false); +initArrays(&Gh, N, false, false); +initArrays(&Ad, N, true, false); +initArrays(&Bd, N, true, false); +initArrays(&Cd, N, true, false); +initArrays(&Dd, N, true, false); - setArray(Dh, N, T(1)); - setArray(Eh, N, T(1)); - setArray(Bh, N, T(1)); +setArray(Dh, N, T(1)); +setArray(Eh, N, T(1)); +setArray(Bh, N, T(1)); - H2D(Ad, Dh, size); - H2D(Bd, Eh, size); +H2D(Ad, Dh, size); +H2D(Bd, Eh, size); - D2HAsync(Ah, Ad, size, stream); - D2DAsync(Cd, Bd, size, stream); - H2HAsync(Ch, Bh, size, stream); - H2DAsync(Dd, Ch, size, stream); - hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Dd); +D2HAsync(Ah, Ad, size, stream); +D2DAsync(Cd, Bd, size, stream); +H2HAsync(Ch, Bh, size, stream); +H2DAsync(Dd, Ch, size, stream); +hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Dd); - D2H(Fh, Cd, size); - D2H(Gh, Dd, size); +D2H(Fh, Cd, size); +D2H(Gh, Dd, size); - HIPCHECK(hipDeviceSynchronize()); - HIPASSERT(Dh[10] == Ah[10]); - HIPASSERT(Eh[10] == Fh[10]); - HIPASSERT(Bh[10] + T(1) == Gh[10]); -} - -int main(int argc, char *argv[]) -{ - HipTest::parseStandardArguments(argc, argv, true); - - test12345(); - test13452(); - test14523(); - test15234(); - - test23451(); - test24513(); - test25134(); - test21345(); - - test34512(); - test35124(); - test31245(); - test32451(); - - test45123(); - test41235(); - test42351(); - test43512(); - - test51234(); - test52341(); - test53412(); - test54123(); - - passed(); +HIPCHECK(hipDeviceSynchronize()); +HIPASSERT(Dh[10] == Ah[10]); +HIPASSERT(Eh[10] == Fh[10]); +HIPASSERT(Bh[10] + T(1) == Gh[10]); } +int main(){ +test12345(); +test13452(); +test14523(); +test15234(); + +test23451(); +test24513(); +test25134(); +test21345(); + +test34512(); +test35124(); +test31245(); +test32451(); + +test45123(); +test41235(); +test42351(); +test43512(); + +test51234(); +test52341(); +test53412(); +test54123(); +passed(); + +}