SWDEV-273943 - Enable runtimeApi.event dtest

Add device_id_ in hip::event to match cuda behaviour in
hipEventQuery() and hipEventRecord().
Enable hipEventElapsedTime test on AMD platform.
Workarround sporadic crash of hipEventIpc test due to
some bug of event ipc.
Add missing hipEventDestroy() in some event tests.
Fix some logic code errors.
Fix typo in comment.

Change-Id: I9ec74c475161b3e31df48d193449023e921f2924

hipamd/tests/src/runtimeApi/event/hipEvent.cpp

@@ -68,6 +68,7 @@ bool hipEvent_Nullcheck() {
     printf("hipEventDestroy failed when nullptr is passed \n");
     TestStatus = false;
   }
+
   return TestStatus;
 }
 

hipamd/tests/src/runtimeApi/event/hipEventElapsedTime.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
  * TEST: %t
  * HIT_END
  */
@@ -49,6 +49,8 @@ void NegativeTests(){
         HIPCHECK(hipEventCreateWithFlags(&start,hipEventDisableTiming));
         HIPCHECK(hipEventCreateWithFlags(&stop,hipEventDisableTiming));
         HIPASSERT(hipEventElapsedTime(&timeElapsed, start, stop) == hipErrorInvalidHandle);
+        HIPCHECK(hipEventDestroy(start));
+        HIPCHECK(hipEventDestroy(stop));
     }
 
     // events created different devices
@@ -59,7 +61,9 @@ void NegativeTests(){
             // create event on dev=0
             HIPCHECK(hipSetDevice(0));
             hipEvent_t start;
+            hipEvent_t start1;
             HIPCHECK(hipEventCreate(&start));
+            HIPCHECK(hipEventCreate(&start1));
 
             HIPCHECK(hipEventRecord(start, nullptr));
             HIPCHECK(hipEventSynchronize(start));
@@ -69,11 +73,19 @@ void NegativeTests(){
             hipEvent_t stop;
             HIPCHECK(hipEventCreate(&stop));
 
+            // start1 on device 0 but null stream on device 1
+            HIPASSERT(hipEventRecord(start1, nullptr) == hipErrorInvalidHandle);
+
             HIPCHECK(hipEventRecord(stop, nullptr));
             HIPCHECK(hipEventSynchronize(stop));
 
             float tElapsed = 1.0f;
+            // start on device 0 but stop on device 1
             HIPASSERT(hipEventElapsedTime(&tElapsed,start,stop) == hipErrorInvalidHandle);
+
+            HIPCHECK(hipEventDestroy(start));
+            HIPCHECK(hipEventDestroy(start1));
+            HIPCHECK(hipEventDestroy(stop));
         }
     }
 }
@@ -92,7 +104,10 @@ void PositiveTest(){
     HIPCHECK(hipEventSynchronize(stop));
 
     float tElapsed = 1.0f;
-    HIPCHECK(hipEventElapsedTime(&tElapsed,start,stop));
+    HIPCHECK(hipEventElapsedTime(&tElapsed, start, stop));
+
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
 }
 
 int main(){

hipamd/tests/src/runtimeApi/event/hipEventIpc.cpp

@@ -78,7 +78,7 @@ int main(int argc, char* argv[]) {
 
 
         float eventMs = 1.0f;
-        // should fail
+        // should fail due to hipEventDisableTiming
         HIPASSERT(hipSuccess != hipEventElapsedTime(&eventMs, start, stop));
         float hostMs = HipTest::elapsed_time(hostStart, hostStop);
 
@@ -97,16 +97,27 @@ int main(int argc, char* argv[]) {
     HIPCHECK(hipEventSynchronize(ipc_event));
 
     HIPCHECK(hipEventDestroy(ipc_event));
+
+#ifndef __HIP_PLATFORM_AMD__
     HIPCHECK(hipEventDestroy(start));
     HIPCHECK(hipEventDestroy(stop));
+#endif
 
     HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
-
     printf("check:\n");
 
+#ifdef __HIP_PLATFORM_AMD__
+    // Due to implementation bug of hipEventInterprocess, as a workaround,
+    // we have to move hipEventDestroy here. Otherwise sporadic crash will
+    // happen in above hipMemcpy(). If hipEventInterprocess is officially
+    // implemented, we should revisit here and move these back to match cuda
+    // event behavior.
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
+#endif
+
     HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
 
-
     passed();
 }

hipamd/tests/src/runtimeApi/event/hipEventRecord.cpp

@@ -91,6 +91,8 @@ int main(int argc, char* argv[]) {
 
     HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
 
     printf("check:\n");
 

hipamd/tests/src/runtimeApi/event/record_event.cpp

@@ -104,11 +104,10 @@ void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_
             assert(0);
     };
 
-
     float t;
 
     hipError_t e = hipEventElapsedTime(&t, start, start);
-    if ((e != hipSuccess) && (e != hipErrorNotReady)) {
+    if ((e != hipSuccess) && (e != hipErrorNotReady || syncMode != syncNone)) {
         failed("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e));
     }
 
@@ -143,6 +142,9 @@ void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_
         HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidHandle);
     }
 
+    HIPCHECK(hipEventDestroy(neverRecorded));
+    HIPCHECK(hipEventDestroy(timingDisabled));
+
     HIPCHECK(hipEventDestroy(start));
     HIPCHECK(hipEventDestroy(stop));