diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 84a19815be..d47728759f 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -8,6 +8,8 @@ include_directories( ${PROJECT_SOURCE_DIR}/include ) set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) +set (HIP_BUILD_LOCAL 0) + set(HIP_PATH $ENV{HIP_PATH}) if (NOT DEFINED HIP_PATH) set (HIP_PATH ../..) @@ -42,8 +44,11 @@ if (${HIP_PLATFORM} STREQUAL "hcc") # This will create a subdir "hip_hcc" in the test build directory # Any changes to hip_hcc source will be detected and force the library and then the tests to be rebuilt. - add_subdirectory(${HIP_PATH} build.hip_hcc) - link_directories(${CMAKE_CURRENT_BINARY_DIR}/build.hip_hcc) # search the local hip_hcc for libhip_hcc.a + if (${HIP_BUILD_LOCAL}) + add_subdirectory(${HIP_PATH} build.hip_hcc) + #link_directories(${CMAKE_CURRENT_BINARY_DIR}/build.hip_hcc) # search the local hip_hcc for libhip_hcc.a + set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) + endif() elseif (${HIP_PLATFORM} STREQUAL "nvcc") @@ -62,7 +67,6 @@ endif() set (HIPCC ${HIP_PATH}/bin/hipcc) set (CMAKE_CXX_COMPILER ${HIPCC}) -#set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) add_library(test_common OBJECT test_common.cpp ) @@ -72,7 +76,9 @@ add_library(test_common OBJECT test_common.cpp ) macro (make_hip_executable exe cpp) if (${HIP_PLATFORM} STREQUAL "hcc") add_executable (${exe} ${cpp} ${ARGN} $ ) - target_link_libraries(${exe} hip_hcc) + if (${HIP_BUILD_LOCAL}) + target_link_libraries(${exe} hip_hcc) + endif() else() add_executable (${exe} ${cpp} ${ARGN} $ ) endif() @@ -142,7 +148,7 @@ make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisio make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp) #TODO - re-enable. This uses the pointer add feature. -#make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) +make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp) make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) make_hip_executable (hipHostAlloc hipHostAlloc.cpp) @@ -173,8 +179,9 @@ make_test(hipGridLaunch " " ) make_test(hipEnvVarDriver " " ) #TODO -reenable #make_test(hipPointerAttrib " " ) -#make_test(hipMultiThreadStreams1 " " ) -#make_test(hipMultiThreadStreams2 " " ) +#make_test(hipMultiThreadStreams1 " " ) Fails if 0x3 specified, passes otherwise. + +make_test(hipMultiThreadStreams2 " " ) make_test(hipMemcpy_simple " " ) make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 ) make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 ) @@ -198,6 +205,8 @@ make_test(hipFuncSetDeviceFlags " ") make_test(hipFuncGetDevice " ") make_test(hipFuncSetDevice " ") make_test(hipFuncDeviceSynchronize " ") -make_test (hipThreadSafeDevice " ") +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-serial" --tests 0x1) +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-pyramid" --tests 0x4) +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-nearzero" --tests 0x10) make_hipify_test(specialFunc.cu ) diff --git a/tests/src/hipMultiThreadStreams1.cpp b/tests/src/hipMultiThreadStreams1.cpp index a3dd94e077..2dc8cb27e7 100644 --- a/tests/src/hipMultiThreadStreams1.cpp +++ b/tests/src/hipMultiThreadStreams1.cpp @@ -84,11 +84,11 @@ void test_multiThread_1(std::string testName, hipStream_t stream0, hipStream_t s std::cout << testName << std::endl; // Test 2 threads operating on same stream: - std::thread t1 (simpleVectorCopy, 2000000/*mb*/, 1000, stream0); + std::thread t1 (simpleVectorCopy, 2000000/*mb*/, 100/*iters*/, stream0); if (serialize) { t1.join(); } - std::thread t2 (simpleVectorCopy, 2000000/*mb*/, 1000, stream1); + std::thread t2 (simpleVectorCopy, 2000000/*mb*/, 100/*iters*/, stream1); if (serialize) { t2.join(); } @@ -119,19 +119,21 @@ int main(int argc, char *argv[]) simpleVectorCopy (2000000/*mb*/, 10/*iters*/, stream); simpleVectorCopy (2000000/*mb*/, 10/*iters*/, stream); - //HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipStreamDestroy(stream)); } - if (p_tests & 0x2) { - hipStream_t stream0, stream1; - HIPCHECK (hipStreamCreate(&stream0)); - HIPCHECK (hipStreamCreate(&stream1)); + hipStream_t stream0, stream1; + HIPCHECK (hipStreamCreate(&stream0)); + HIPCHECK (hipStreamCreate(&stream1)); + if (p_tests & 0x2) { // Easy tests to verify the test works - these don't allow overlap between the threads: test_multiThread_1 ("Multithread NULL with serialized", NULL, NULL, true); - test_multiThread_1 ("Multithread with serialized", stream0, stream1, true); + test_multiThread_1 ("Multithread two streams serialized", stream0, stream1, true); + } + if (p_tests & 0x4) { test_multiThread_1 ("Multithread with NULL stream", NULL, NULL, false); test_multiThread_1 ("Multithread with two streams", stream0, stream1, false); test_multiThread_1 ("Multithread with one stream", stream0, stream0, false); diff --git a/tests/src/hipThreadSafeDevice.cpp b/tests/src/hipThreadSafeDevice.cpp index 0534d6fbae..a1f64aceb3 100644 --- a/tests/src/hipThreadSafeDevice.cpp +++ b/tests/src/hipThreadSafeDevice.cpp @@ -74,7 +74,7 @@ void multiThread_pyramid(bool serialize, int iters) // Create 3 streams, all creating and destroying streams on the same device. // Try to keep number of streams near zero, to cause problems. -void multiThread_tiny(bool serialize, int iters) +void multiThread_nearzero(bool serialize, int iters) { printf ("%s creating %d streams x 3 threads\n", __func__, iters); std::thread t1 (createThenDestroyStreams, iters, 1); @@ -96,9 +96,9 @@ void multiThread_tiny(bool serialize, int iters) } if (!serialize) { - t1.join(); - t2.join(); - t3.join(); + t1.join(); printf ("t1 done\n"); + t2.join(); printf ("t2 done\n"); + t3.join(); printf ("t3 done\n"); } } @@ -113,7 +113,8 @@ int main(int argc, char *argv[]) createThenDestroyStreams(10, 10); }; - if (p_tests & 0x2) { + /*disable, this takess a while and if the next one works then no need to run serial*/ + if (1 && (p_tests & 0x2)) { printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n"); multiThread_pyramid(true, 10); } @@ -129,8 +130,8 @@ int main(int argc, char *argv[]) // } if (p_tests & 0x10) { - printf ("\ntest 0x10 : parallel multiThread_tiny(1000) \n"); - multiThread_tiny(false, 1000); + printf ("\ntest 0x10 : parallel multiThread_nearzero(1000) \n"); + multiThread_nearzero(false, 1000); } passed(); diff --git a/tests/src/test_common.h b/tests/src/test_common.h index d92692c978..ce85b3898a 100644 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include +#include #include #include @@ -232,7 +233,10 @@ void checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch=true } mismatchCount++; if ((mismatchCount <= mismatchesToPrint) && expectMatch) { - std::cout << "At " << i << " Computed:" << result_H[i] << ", expected:" << expected << std::endl; + std::cout << std::fixed << std::setprecision(32); + std::cout << "At " << i << std::endl; + std::cout << " Computed:" << result_H[i] << std::endl; + std::cout << " Expected:" << expected << std::endl; } } }