diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 44becb0ed3..e1042a2d54 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -69,10 +69,7 @@ macro (make_hip_executable exe cpp) endif() endmacro() - -macro (make_test exe ) - string (REPLACE " " "" smush_args ${ARGN}) - set (testname ${exe}${smush_args}.tst) +macro (make_named_test exe testname ) add_test (NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} ) @@ -81,6 +78,13 @@ macro (make_test exe ) ) endmacro() +macro (make_test exe ) + string (REPLACE " " "" smush_args ${ARGN}) + set (testname ${exe}${smush_args}.tst) + + make_named_test(${exe} ${testname} ${ARGN}) +endmacro() + macro (make_hipify_test sourceFile ) #string (REPLACE " " "" smush_args ${ARGN}) @@ -112,6 +116,7 @@ make_hip_executable (hip_ffs hip_ffs.cpp) make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) make_hip_executable (hipEnvVar hipEnvVar.cpp) make_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) +make_hip_executable (hipMemcpy_simple hipMemcpy_simple.cpp) make_hip_executable (hipMemcpy hipMemcpy.cpp) make_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp) make_hip_executable (hipMemset hipMemset.cpp) @@ -147,7 +152,10 @@ make_test(hipPointerAttrib " " ) make_test(hipMultiThreadStreams1 " " ) make_test(hipMultiThreadStreams2 " " ) -make_test(hipMemcpy " " ) +make_test(hipMemcpy_simple " " ) +make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 ) +make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 ) +make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 ) make_test(hipMemcpyAsync " " ) make_test(hipHcc " " ) diff --git a/projects/hip/tests/src/hipMemcpy.cpp b/projects/hip/tests/src/hipMemcpy.cpp index b2bfc63fe5..7dc8de8d2e 100644 --- a/projects/hip/tests/src/hipMemcpy.cpp +++ b/projects/hip/tests/src/hipMemcpy.cpp @@ -28,39 +28,6 @@ void printSep() printf ("======================================================================================\n"); } -//--- -// Test simple H2D copies and back. -// Designed to stress a small number of simple smoke tests -void simpleTest1() -{ - printf ("test: %s\n", __func__); - size_t Nbytes = N*sizeof(int); - printf ("N=%zu Nbytes=%6.2fMB\n", N, Nbytes/1024.0/1024.0); - - int *A_d, *B_d, *C_d; - int *A_h, *B_h, *C_h; - - HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - - printf ("A_d=%p B_d=%p C_d=%p A_h=%p B_h=%p C_h=%p\n", A_d, B_d, C_d, A_h, B_d, C_h); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); - - 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); - - HIPCHECK ( hipMemcpy(C_h, C_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 (hipDeviceReset()); - - printf (" %s success\n", __func__); -} @@ -148,7 +115,7 @@ void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, boo //--- //Try all the 16 possible combinations to memcpytest2 - usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault template -void memcpytest2_loop(size_t numElements) +void memcpytest2_for_type(size_t numElements) { printSep(); @@ -220,6 +187,7 @@ void multiThread_1(bool serialize, bool usePinnedHost) + int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true); @@ -229,19 +197,37 @@ int main(int argc, char *argv[]) if (p_tests & 0x1) { + printf ("\n\n=== tests&1 (types and different memcpy kinds (H2D, D2H, H2H, D2D)\n"); HIPCHECK ( hipDeviceReset() ); - simpleTest1(); + memcpytest2_for_type(N); + memcpytest2_for_type(N); + memcpytest2_for_type(N); + memcpytest2_for_type(N); + printf ("===\n\n\n"); } + if (p_tests & 0x2) { - HIPCHECK ( hipDeviceReset() ); - memcpytest2_loop(N); - memcpytest2_loop(N); - memcpytest2_loop(N); - memcpytest2_loop(N); + // Some tests around the 64MB boundary which have historically shown issues: + printf ("\n\n=== tests&0x2 (64MB boundary)\n"); +#if 0 + // These all pass: + memcpytest2(15*1024*1024, 1, 0, 0, 0); + memcpytest2(16*1024*1024, 1, 0, 0, 0); + memcpytest2(16*1024*1024+16*1024, 1, 0, 0, 0); +#endif + // Just over 64MB: + memcpytest2(16*1024*1024+512*1024, 1, 0, 0, 0); + memcpytest2(17*1024*1024+1024, 1, 0, 0, 0); + memcpytest2(32*1024*1024, 1, 0, 0, 0); + memcpytest2(32*1024*1024, 0, 0, 0, 0); + memcpytest2(32*1024*1024, 1, 1, 1, 0); + memcpytest2(32*1024*1024, 1, 1, 1, 0); } + if (p_tests & 0x4) { + printf ("\n\n=== tests&4 (test sizes and offsets)\n"); HIPCHECK ( hipDeviceReset() ); printSep(); memcpytest2_sizes(0,0); @@ -254,6 +240,7 @@ int main(int argc, char *argv[]) } if (p_tests & 0x8) { + printf ("\n\n=== tests&8\n"); HIPCHECK ( hipDeviceReset() ); printSep(); @@ -269,8 +256,10 @@ int main(int argc, char *argv[]) // Remove serialization, and use unpinned. multiThread_1(false, false); // TODO + printf ("===\n\n\n"); } + passed(); } diff --git a/projects/hip/tests/src/hipMemcpy_simple.cpp b/projects/hip/tests/src/hipMemcpy_simple.cpp new file mode 100644 index 0000000000..30220880c9 --- /dev/null +++ b/projects/hip/tests/src/hipMemcpy_simple.cpp @@ -0,0 +1,132 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include "hip_runtime.h" +#include "test_common.h" + + +//--- +// Test simple H2D copies and back. +// Designed to stress a small number of simple smoke tests +void simpleTest1() +{ + printf ("test: %s\n", __func__); + size_t Nbytes = N*sizeof(int); + printf ("N=%zu Nbytes=%6.2fMB\n", N, Nbytes/1024.0/1024.0); + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + printf ("A_d=%p B_d=%p C_d=%p A_h=%p B_h=%p C_h=%p\n", A_d, B_d, C_d, A_h, B_d, C_h); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + 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); + + HIPCHECK ( hipMemcpy(C_h, C_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 (hipDeviceReset()); + + printf (" %s success\n", __func__); +} + + +template +void simpleTest2(size_t numElements, bool usePinnedHost) +{ + printf ("test: %s<%s>\n", __func__, TYPENAME(T)); + size_t sizeElements = numElements * sizeof(T); + + T *A_d, *A_h1, *A_h2; + + if (usePinnedHost) { + HIPCHECK ( hipMallocHost(&A_h1, sizeElements) ); + HIPCHECK ( hipMallocHost(&A_h2, sizeElements) ); + } else { + A_h1 = (T*)malloc(sizeElements); + A_h2 = (T*)malloc(sizeElements); + } + + // Alloc device array: + HIPCHECK ( hipMalloc(&A_d, sizeElements) ); + + + for (size_t i=0; i(N, true/*usePinnedHost*/); + simpleTest2(N, true/*usePinnedHost*/); + } + + if (p_tests & 0x4) { + printf ("\n\n=== tests&2 (copy pin-pong, unpinned host)\n"); + simpleTest2(N, false/*usePinnedHost*/); + simpleTest2(N, false/*usePinnedHost*/); + } + + passed(); +};