AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Test improvements

Telusuri Sumber 
- partition hipThreadSafeDevice into smaller pieces.
- Add debug to hipMultiThreadStream.
- print more precision when mismatch detected.
- enable more tests in CMakeFiles.txt.


[ROCm/clr commit: 98a766ba75]
This commit is contained in:
Ben Sander
2016-03-29 17:29:31 -05:00
orang tua f1295bc03d
melakukan 875a8a2c84
4 mengubah file dengan 40 tambahan dan 24 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/tests/src/CMakeLists.txt
+17 -8
Melihat File
@@ -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_OBJECTS:test_common> )
target_link_libraries(${exe} hip_hcc)
if (${HIP_BUILD_LOCAL})
target_link_libraries(${exe} hip_hcc)
endif()
else()
add_executable (${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common> )
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 )
projects/clr/hipamd/tests/src/hipMultiThreadStreams1.cpp
+10 -8
Melihat File
@@ -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<T, HipTest::Pinned, C>, 2000000/*mb*/, 1000, stream0);
std::thread t1 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream0);
if (serialize) {
t1.join();
}
std::thread t2 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 1000, stream1);
std::thread t2 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream1);
if (serialize) {
t2.join();
}
@@ -119,19 +119,21 @@ int main(int argc, char *argv[])
simpleVectorCopy<float, HipTest::Pinned, HipTest::MemcpyAsync> (2000000/*mb*/, 10/*iters*/, stream);
simpleVectorCopy<float, HipTest::Pinned, HipTest::Memcpy> (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<float, HipTest::MemcpyAsync> ("Multithread NULL with serialized", NULL, NULL, true);
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with serialized", stream0, stream1, true);
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread two streams serialized", stream0, stream1, true);
}
if (p_tests & 0x4) {
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with NULL stream", NULL, NULL, false);
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with two streams", stream0, stream1, false);
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with one stream", stream0, stream0, false);
projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp
+8 -7
Melihat File
@@ -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();
projects/clr/hipamd/tests/src/test_common.h
+5 -1
Melihat File
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <iostream>
#include <iomanip>
#include <sys/time.h>
#include <stddef.h>
@@ -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;
}
}
}