Update tests.

Fix some NVCC issues.
Add hipStreamSync2, record_event tests.

tests/src/runtimeApi/stream/hipStreamSync2.cpp

@@ -0,0 +1,169 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+
+#include "test_common.h"
+
+enum SyncMode {
+    syncNone,
+    syncNullStream,
+    syncOtherStream,
+    syncMarkerThenOtherStream,
+    syncMarkerThenOtherNonBlockingStream,
+    syncDevice
+};
+
+
+const char *syncModeString(int syncMode) {
+    switch (syncMode) {
+        case syncNone:
+            return "syncNone";
+        case syncNullStream:
+            return "syncNullStream";
+        case syncOtherStream:
+            return "syncOtherStream";
+        case syncMarkerThenOtherStream:
+            return "syncMarkerThenOtherStream";
+        case syncMarkerThenOtherNonBlockingStream:
+            return "syncMarkerThenOtherNonBlockingStream";
+        case syncDevice:
+            return "syncDevice";
+        default:
+            return "unknown";
+    };
+};
+
+
+void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expectMismatch)
+{
+    printf ("\ntest: syncMode=%s expectMismatch=%d\n", syncModeString(syncMode), expectMismatch);
+
+    size_t sizeBytes = numElements * sizeof(int);
+
+    int count =100;
+    int init0 = 0;
+    HIPCHECK(hipMemset(C_d, init0, sizeBytes));
+    for (int i=0; i<numElements; i++) {
+        C_h[i] = -1; // initialize
+    }
+
+    hipStream_t otherStream = 0;
+    unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ?  hipStreamNonBlocking : hipStreamDefault;
+    HIPCHECK(hipStreamCreateWithFlags(&otherStream, flags));
+    hipEvent_t e;
+    HIPCHECK(hipEventCreate(&e));
+
+
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
+    // Launch kernel into null stream, should result in C_h == count.
+    hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, 0 /*stream*/,    C_d, C_h, numElements, count);
+
+    switch (syncMode) {
+        case syncNone:
+            break;
+        case syncNullStream:
+            HIPCHECK(hipStreamSynchronize(0));  // wait on host for null stream:
+            break;
+        case syncOtherStream:
+            // Does this synchronize with the null stream?
+            HIPCHECK(hipStreamSynchronize(otherStream));  
+            break;
+        case syncMarkerThenOtherStream:
+        case syncMarkerThenOtherNonBlockingStream:
+            HIPCHECK(hipEventRecord(e, otherStream));  // this may wait for NULL stream depending hipStreamNonBlocking flag above
+            HIPCHECK(hipStreamSynchronize(otherStream));  
+            break;
+        case syncDevice:
+            HIPCHECK(hipDeviceSynchronize());  
+            break;
+        default:
+            assert(0);
+    };
+
+    int mismatches = 0;
+    int expected = init0 + count;
+    for (int i=0; i<numElements; i++) {
+        bool compareEqual = (C_h[i] == expected);
+        if (!compareEqual) {
+            mismatches ++;
+            if  (!expectMismatch) {
+                printf ("C_h[%d] (%d) != %d\n", i, C_h[i], expected);
+                assert(C_h[i] == expected);
+            }
+        }
+    }
+
+    if (expectMismatch) {
+        assert (mismatches > 0);
+    }
+
+
+    HIPCHECK(hipStreamDestroy(otherStream));
+    HIPCHECK(hipEventDestroy(e));
+
+    printf ("test:   OK - %d mismatches (%6.2f%%)\n",  mismatches, ((double)(mismatches)*100.0)/numElements);
+}
+
+
+void testEventRecord()
+{
+}
+
+
+void runTests(int64_t numElements)
+{
+    size_t sizeBytes = numElements * sizeof(int);
+
+    printf ("\n\ntest: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes, sizeBytes/1024.0/1024.0);
+
+
+    int *C_h, *C_d;
+    HIPCHECK(hipMalloc(&C_d, sizeBytes));
+    HIPCHECK(hipHostMalloc(&C_h, sizeBytes));
+
+
+    {
+        test (C_d, C_h, numElements,  syncNone,                             true /*expectMismatch*/);
+        test (C_d, C_h, numElements,  syncNullStream,                       false /*expectMismatch*/);
+        test (C_d, C_h, numElements,  syncOtherStream,                      true /*expectMismatch*/);
+        test (C_d, C_h, numElements,  syncDevice,                           false /*expectMismatch*/);
+        test (C_d, C_h, numElements,  syncMarkerThenOtherStream,            false /*expectMismatch*/);
+        test (C_d, C_h, numElements,  syncMarkerThenOtherNonBlockingStream, true /*expectMismatch*/);
+    }
+
+
+    HIPCHECK(hipFree(C_d));
+    HIPCHECK(hipHostFree(C_h));
+}
+
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/);
+
+    runTests(40000000);
+
+    passed();
+}