Cleanup hipEvent. (Intermediate checkpoint)

Support hipEventDisableSystemRelease flag.
Update test.
Remove stray printf

src/hip_event.cpp

@@ -82,7 +82,7 @@ void ihipEvent_t::setTimestamp()
     }
 
     if (_state != hipEventStatusRecorded) {
-        printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
+        //printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
     }
 }
 
@@ -92,7 +92,10 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
     hipError_t e = hipSuccess;
 
     // TODO-IPC - support hipEventInterprocess.
-    unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming;
+    unsigned supportedFlags =   hipEventDefault 
+                              | hipEventBlockingSync 
+                              | hipEventDisableTiming 
+                              | hipEventDisableSystemRelease;
     if ((flags & ~supportedFlags) == 0) {
         ihipEvent_t *eh = new ihipEvent_t(flags);
 
@@ -197,20 +200,18 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
 {
     HIP_INIT_API(ms, start, stop);
 
-    ihipEvent_t *start_eh = start;
-    ihipEvent_t *stop_eh = stop;
-
     start->setTimestamp();
     stop->setTimestamp();
 
     hipError_t status = hipSuccess;
     *ms = 0.0f;
 
-    if (start_eh && stop_eh) {
-        if ((start_eh->_state == hipEventStatusRecorded) && (stop_eh->_state == hipEventStatusRecorded)) {
+    if (start && stop) {
+        // refresh status:
+        if ((start->_state == hipEventStatusRecorded) && (stop->_state == hipEventStatusRecorded)) {
             // Common case, we have good information for both events.
 
-            int64_t tickDiff = (stop_eh->timestamp() - start_eh->timestamp());
+            int64_t tickDiff = (stop->timestamp() - start->timestamp());
 
             uint64_t freqHz;
             hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
@@ -223,13 +224,16 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
             }
 
 
-        } else if ((start_eh->_state == hipEventStatusRecording) ||
-                   (stop_eh->_state  == hipEventStatusRecording)) {
+        } else if ((start->_state == hipEventStatusRecording) ||
+                   (stop->_state  == hipEventStatusRecording)) {
+
             status = hipErrorNotReady;
-        } else if ((start_eh->_state == hipEventStatusUnitialized) ||
-                   (stop_eh->_state  == hipEventStatusUnitialized)) {
+        } else if ((start->_state == hipEventStatusUnitialized) ||
+                   (stop->_state  == hipEventStatusUnitialized)) {
             status = hipErrorInvalidResourceHandle;
         }
+    } else {
+        status = hipErrorInvalidResourceHandle;
     }
 
     return ihipLogStatus(status);

src/hip_hcc.cpp

@@ -49,7 +49,7 @@ THE SOFTWARE.
 
 
 // needs HCC change for hc::no_scope
-#define USE_NO_SCOPE 0
+#define USE_NO_SCOPE 1
 
 //=================================================================================================
 //Global variables:
@@ -331,10 +331,10 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event)
     LockedAccessor_StreamCrit_t crit(_criticalData);
 
 #if USE_NO_SCOPE
-    printf ("create_marker, flags = %x\n", event->_flags);
+    //printf ("create_marker, flags = %x\n", event->_flags);
     event->_marker = crit->_av.create_marker((event->_flags & hipEventDisableSystemRelease) ? hc::no_scope : hc::system_scope);
 #else
-    event->_marker = crit->_av.create_marker();
+    event->_marker = crit->_av.create_marker((event->_flags & hipEventDisableSystemRelease) ? hc::accelerator_scope : hc::system_scope);
 #endif
 };
 

tests/src/runtimeApi/memory/hipHostMalloc.cpp

@@ -42,6 +42,63 @@ __global__ void Set(int *Ad, int val){
     Ad[tx] = val;
 }
 
+
+#define SYNC_EVENT 0
+#define SYNC_STREAM 1
+#define SYNC_DEVICE 2
+
+std::vector<std::string> syncMsg = {"event", "stream", "device"};
+
+void CheckHostPointer(int numElements, int *ptr, int syncMethod, std::string msg)
+{
+    std::cerr << "test: CheckHostPointer " << msg 
+              << " ptr=" << ptr 
+              << " syncMethod=" << syncMsg[syncMethod] << "\n";
+
+    hipStream_t s;
+    hipEvent_t e;
+
+    // Init:
+    HIPCHECK(hipStreamCreate(&s));
+    HIPCHECK(hipEventCreateWithFlags(&e, hipEventDisableSystemRelease));
+    dim3 dimBlock(64,1,1);
+    dim3 dimGrid(numElements/dimBlock.x,1,1);
+
+    const int expected = 13;
+
+    // Init array to know state:
+    hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42);
+    HIPCHECK(hipDeviceSynchronize());
+
+    hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, ptr, expected);
+    HIPCHECK(hipEventRecord(e, s));
+
+    // Host waits for event :
+    switch (syncMethod) {
+        case SYNC_EVENT:
+            HIPCHECK(hipEventSynchronize(e));
+            break;
+        case SYNC_STREAM:
+            HIPCHECK(hipStreamSynchronize(s));
+            break;
+        case SYNC_DEVICE:
+            HIPCHECK(hipDeviceSynchronize());
+            break;
+        default:
+            assert(0);
+    };
+            
+    for (int i=0; i<numElements; i++) {
+        if (ptr[i] != expected) {
+            printf ("mismatch at %d: %d != %d\n", i, ptr[i], expected);
+            assert(ptr[i] == expected);
+        }
+    }
+
+    HIPCHECK(hipStreamDestroy(s));
+    HIPCHECK(hipEventDestroy(e));
+};
+
 int main(){
 
 
@@ -86,38 +143,41 @@ int main(){
     }
 
     {
-        int *A, *B;
         int numElements = 1024*16;
         size_t sizeBytes = numElements * sizeof (int);
-#ifdef __HIP_PLATFORM_HCC__
-        HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
 
-        assert (A == 0);
+#ifdef __HIP_PLATFORM_HCC__
+        { 
+            // Stimulate error condition:
+            int *A = &numElements;
+            HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
+
+            assert (A == 0);
+        }
 #endif
 
-        HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
-        hipStream_t s;
-        hipEvent_t e;
 
-        // Init:
-        HIPCHECK(hipStreamCreate(&s));
-        HIPCHECK(hipEventCreateWithFlags(&e, 0));
-        dim3 dimBlock(64,1,1);
-        dim3 dimGrid(numElements/dimBlock.x,1,1);
+        {
+            int *A = nullptr;
+            HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocNonCoherent));
+            const char *ptrType = "non-coherent"; // TODO
+            //CheckHostPointer(numElements, A, SYNC_DEVICE,   ptrType);
+            //CheckHostPointer(numElements, A, SYNC_STREAM,  ptrType);
+            CheckHostPointer(numElements, A, SYNC_EVENT,   ptrType);
+        }
 
-        // Init array to know state:
-        hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, A, -42);
-        HIPCHECK(hipDeviceSynchronize());
+        if (0) { // TODO, remove me
+            int *A = nullptr;
+            HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
+            const char *ptrType = "coherent";
+            CheckHostPointer(numElements, A, SYNC_DEVICE,   ptrType);
+            CheckHostPointer(numElements, A, SYNC_STREAM,  ptrType);
+            CheckHostPointer(numElements, A, SYNC_EVENT,   ptrType);
+        }
 
-        hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, A, 13);
-        HIPCHECK(hipEventRecord(e, s));
 
-        // Host waits for event :
-        HIPCHECK(hipEventSynchronize(e));
 
-        // check result?
 
-        HIPCHECK(hipHostMalloc((void**)&B, sizeBytes, hipHostMallocNonCoherent));
     }
         
     passed();