SWDEV-294591 - Fix hipEventIpc failure on NV

make hipIpcOpenEventHandle has the same behavour of cudaIpcOpenEventHandle.
Add Api usages.

Change-Id: I4248b2cebd3de156f9d5d427e1797da22fb964eb


[ROCm/hip commit: c053d60282]

projects/hip/include/hip/hip_runtime_api.h

@@ -1274,8 +1274,40 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
  *
  */
 hipError_t hipIpcCloseMemHandle(void* devPtr);
+
+/**
+ * @brief Gets an opaque interprocess handle for an event.
+ *
+ * This opaque handle may be copied into other processes and opened with cudaIpcOpenEventHandle.
+ * Then cudaEventRecord, cudaEventSynchronize, cudaStreamWaitEvent and cudaEventQuery may be used in
+ * either process. Operations on the imported event after the exported event has been freed with hipEventDestroy
+ * will result in undefined behavior.
+ *
+ * @param[out]  handle Pointer to cudaIpcEventHandle to return the opaque event handle
+ * @param[in]   event  Event allocated with cudaEventInterprocess and cudaEventDisableTiming flags
+ *
+ * @returns #hipSuccess, #hipErrorInvalidConfiguration, #hipErrorInvalidValue
+ *
+ */
 hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event);
+
+/**
+ * @brief Opens an interprocess event handles.
+ *
+ * Opens an interprocess event handle exported from another process with cudaIpcGetEventHandle. The returned
+ * hipEvent_t behaves like a locally created event with the hipEventDisableTiming flag specified. This event
+ * need be freed with hipEventDestroy. Operations on the imported event after the exported event has been freed
+ * with hipEventDestroy will result in undefined behavior. If the function is called within the same process where
+ * handle is returned by hipIpcGetEventHandle, it will return hipErrorInvalidContext.
+ *
+ * @param[out]  event  Pointer to hipEvent_t to return the event
+ * @param[in]   handle The opaque interprocess handle to open
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidContext
+ *
+ */
 hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
+
 // end doxygen Device
 /**
  * @}
@@ -2485,6 +2517,21 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
  */
 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
                               hipStream_t stream);
+
+/**
+ *  @brief Returns a global pointer from a module.
+ *  Returns in *dptr and *bytes the pointer and size of the global of name name located in module hmod.
+ *  If no variable of that name exists, it returns hipErrorNotFound. Both parameters dptr and bytes are optional.
+ *  If one of them is NULL, it is ignored and hipSuccess is returned.
+ *
+ *  @param[out]  dptr  Returned global device pointer
+ *  @param[out]  bytes Returned global size in bytes
+ *  @param[in]   hmod  Module to retrieve global from
+ *  @param[in]   name  Name of global to retrieve
+ *
+ *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotFound, #hipErrorInvalidContext
+ *
+ */
 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
     hipModule_t hmod, const char* name);
 hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol);

projects/hip/tests/src/runtimeApi/event/hipEventIpc.cpp

@@ -24,7 +24,7 @@ THE SOFTWARE.
 // forces synchronization : set
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../test_common.cpp
  * TEST: %t --iterations 10
  * HIT_END
  */
@@ -87,16 +87,19 @@ int main(int argc, char* argv[]) {
         printf("\n");
 
     }
-
     hipIpcEventHandle_t ipc_handle;
     HIPCHECK(hipIpcGetEventHandle(&ipc_handle, start));
 
     hipEvent_t ipc_event;
-    HIPCHECK(hipIpcOpenEventHandle(&ipc_event, ipc_handle));
+    auto ret = hipIpcOpenEventHandle(&ipc_event, ipc_handle);
 
-    HIPCHECK(hipEventSynchronize(ipc_event));
+    if (ret == hipErrorInvalidContext) {
+      // tests/src/ipc/hipMultiProcIpcEvent.cpp is the right sample in different process.
+      printf("hipIpcOpenEventHandle() should be called in a different process\n");
+    } else if (ret == hipSuccess) {
+      // To be removed
+    }
 
-    HIPCHECK(hipEventDestroy(ipc_event));
     HIPCHECK(hipEventDestroy(start));
     HIPCHECK(hipEventDestroy(stop));