Make ihipDevice_t thread-safe.

Move critical data into separate class and protect with LockAccessor
wrapper class.

For device, the streams list is the critical data since it is modified when
streams are created or destroyed.   The streams list is accessed in
several places including when synchronizing across all streams on the
device (ie from the default stream).
Other device data is set once by the device cosntructor and is not critical
so

All functions which acquire the LockAccessor now named with "locked_" prefix.

tests/src/CMakeLists.txt

@@ -9,7 +9,6 @@ set (HIP_Unit_Test_VERSION_MAJOR 1)
 set (HIP_Unit_Test_VERSION_MINOR 0)
 
 set(HIP_PATH $ENV{HIP_PATH})
-MESSAGE("HIP_PATH=" ${HIP_PATH})
 if (NOT DEFINED HIP_PATH)
 	set (HIP_PATH ../..)
 endif()
@@ -39,12 +38,13 @@ if (${HIP_PLATFORM} STREQUAL "hcc")
 
 	#These includes are used for all files.
 	#Include HIP and HC since the tests need both of these:
-	#Note below HSA path is surgically included only where necessary.
 	include_directories(${HIP_PATH}/include)
-    include_directories(${HSA_PATH}/include)
 
     # 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
+
     
 elseif (${HIP_PLATFORM} STREQUAL "nvcc")
 	MESSAGE ("HIP_PLATFORM=nvcc")
@@ -156,6 +156,8 @@ make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp)
 make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp)
 make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
 
+make_hip_executable (hipThreadSafeDevice hipThreadSafeDevice.cpp) 
+
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
 make_test(hip_popc " " )
@@ -196,5 +198,6 @@ make_test(hipFuncSetDeviceFlags " ")
 make_test(hipFuncGetDevice " ")
 make_test(hipFuncSetDevice " ")
 make_test(hipFuncDeviceSynchronize " ")
+make_test (hipThreadSafeDevice " ") 
 
 make_hipify_test(specialFunc.cu )

tests/src/hipThreadSafeDevice.cpp

@@ -0,0 +1,137 @@
+#include <hip_runtime_api.h>
+#include "test_common.h"
+
+
+// Create a lot of streams and then destroy 'em.
+void createThenDestroyStreams(int iterations, int burstSize)
+{
+    hipStream_t *streams = new hipStream_t[burstSize];
+
+    for (int i=0; i<iterations; i++) {
+        if (p_verbose & 0x1) {
+          printf ("%s iter=%d, create %d then destroy %d\n", __func__, i, burstSize, burstSize);
+        }
+        for (int j=0; j<burstSize; j++) {
+            if (p_verbose & 0x2) {
+                printf ("  %d.%d streamCreate\n", i, j);
+            }
+            HIPCHECK( hipStreamCreate(&streams[j]));
+        }
+        for (int j=0; j<burstSize; j++) {
+            if (p_verbose & 0x2) {
+                printf ("  %d.%d streamDestroy\n", i, j);
+            }
+            HIPCHECK( hipStreamDestroy(streams[j]));
+        }
+    }
+
+    delete streams;
+}
+
+
+void waitStreams(int iterations)
+{
+    // Repeatedly sync and wait for all streams to complete.
+    // TO make this interesting, the test has other threads repeatedly adding and removing streams to the device.
+    for (int i=0; i<iterations; i++) {
+        HIPCHECK(hipDeviceSynchronize());
+    }
+}
+
+
+// Create 3 streams, all creating and destroying streams on the same device.
+// Some create many queue, some not many.
+//
+void multiThread_pyramid(bool serialize, int iters)
+{
+    printf ("%s creating %d streams\n", __func__, iters*100);
+    std::thread t1 (createThenDestroyStreams, iters*1, 100);
+    if (serialize) {
+        t1.join();
+        printf("t1 done\n");
+    }
+
+    std::thread t2 (createThenDestroyStreams, iters*10, 10);
+    if (serialize) {
+        t2.join();
+        printf("t2 done\n");
+    }
+
+    std::thread t3 (createThenDestroyStreams, iters*100, 1);
+    if (serialize) {
+        t3.join();
+        printf("t3 done\n");
+    }
+
+    if (!serialize) {
+        t1.join();
+        t2.join();
+        t3.join();
+    }
+
+}
+
+
+// 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)
+{
+    printf ("%s creating %d streams x 3 threads\n", __func__, iters);
+    std::thread t1 (createThenDestroyStreams, iters, 1);
+    if (serialize) {
+        t1.join();
+        printf("t1 done\n");
+    }
+
+    std::thread t2 (createThenDestroyStreams, iters, 1);
+    if (serialize) {
+        t2.join();
+        printf("t2 done\n");
+    }
+
+    std::thread t3 (waitStreams, iters*50);
+    if (serialize) {
+        t3.join();
+        printf("t3 done\n");
+    }
+
+    if (!serialize) {
+        t1.join();
+        t2.join();
+        t3.join();
+    }
+
+}
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+
+    // Serial version, just call once:
+    if (p_tests & 0x1) {
+        printf ("test 0x1 : serial createThenDestroyStreams(10) \n");
+        createThenDestroyStreams(10, 10);
+    };
+
+    if (p_tests & 0x2) {
+        printf ("test 0x2 : serialized multiThread_1(1) \n");
+        multiThread_pyramid(true, 10);
+    }
+
+    if (p_tests & 0x4) {
+        printf ("test 0x4 : multiThread_pyramid(1) \n");
+        multiThread_pyramid(false, 10);
+    }
+
+    //if (p_tests & 0x8) {
+    //    printf ("test 0x8 : multiThread_pyramid(100) \n");
+    //    multiThread_pyramid(false, 100);
+   // }
+
+    if (p_tests & 0x10) {
+        printf ("test 0x10 : multiThread_tiny(1000) \n");
+        multiThread_tiny(false, 1000);
+    }
+
+    passed();
+}