hip: Add GstHipEvent object

hip event handle wrapper object

Part-of: <https://gitlab.freedesktop.org/gstreamer/gstreamer/-/merge_requests/9281>

subprojects/gst-plugins-bad/sys/hip/gsthip.h

@@ -31,4 +31,5 @@
 #include "gsthiploader.h"
 #include "gsthip-interop.h"
 #include "gsthipstream.h"
+#include "gsthipevent.h"
 

subprojects/gst-plugins-bad/sys/hip/gsthip_fwd.h

@@ -46,5 +46,10 @@ typedef struct _GstHipGraphicsResource GstHipGraphicsResource;
 
 typedef struct _GstHipStream GstHipStream;
 
+typedef struct _GstHipEvent GstHipEvent;
+typedef struct _GstHipEventPool GstHipEventPool;
+typedef struct _GstHipEventPoolClass GstHipEventPoolClass;
+typedef struct _GstHipEventPoolPrivate GstHipEventPoolPrivate;
+
 G_END_DECLS
 

subprojects/gst-plugins-bad/sys/hip/gsthipevent.cpp

@@ -0,0 +1,282 @@
+/* GStreamer
+ * Copyright (C) 2025 Seungha Yang <seungha@centricular.com>
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Library General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Library General Public License for more details.
+ *
+ * You should have received a copy of the GNU Library General Public
+ * License along with this library; if not, write to the
+ * Free Software Foundation, Inc., 51 Franklin St, Fifth Floor,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#ifdef HAVE_CONFIG_H
+#include "config.h"
+#endif
+
+#include "gsthip-config.h"
+#include "gsthip.h"
+#include <mutex>
+#include <queue>
+
+#ifndef GST_DISABLE_GST_DEBUG
+#define GST_CAT_DEFAULT ensure_debug_category()
+static GstDebugCategory *
+ensure_debug_category (void)
+{
+  static GstDebugCategory *cat = nullptr;
+  static std::once_flag once;
+
+  std::call_once (once,[&] {
+        cat = _gst_debug_category_new ("hipevent", 0, "hipevent");
+      });
+
+  return cat;
+}
+#endif
+
+/* *INDENT-OFF* */
+struct _GstHipEvent : public GstMiniObject
+{
+  ~_GstHipEvent ()
+  {
+    if (handle) {
+      auto hip_ret = HipSetDevice (vendor, device_id);
+      if (gst_hip_result (hip_ret, vendor)) {
+        HipEventSynchronize (vendor, handle);
+        HipEventDestroy (vendor, handle);
+      }
+    }
+  }
+
+  GstHipEventPool *pool = nullptr;
+  hipEvent_t handle = nullptr;
+  GstHipVendor vendor;
+  guint device_id;
+};
+
+struct _GstHipEventPoolPrivate
+{
+  ~_GstHipEventPoolPrivate ()
+  {
+    while (!event_pool.empty ()) {
+      auto event = event_pool.front ();
+      event_pool.pop ();
+      gst_mini_object_unref (event);
+    }
+  }
+
+  GstHipVendor vendor;
+  guint device_id;
+  std::mutex lock;
+  std::queue<GstHipEvent *>event_pool;
+};
+/* *INDENT-ON* */
+
+GST_DEFINE_MINI_OBJECT_TYPE (GstHipEvent, gst_hip_event);
+
+static void gst_hip_event_pool_finalize (GObject * object);
+
+#define gst_hip_event_pool_parent_class parent_class
+G_DEFINE_TYPE (GstHipEventPool, gst_hip_event_pool, GST_TYPE_OBJECT);
+
+static void
+gst_hip_event_pool_class_init (GstHipEventPoolClass * klass)
+{
+  auto object_class = G_OBJECT_CLASS (klass);
+
+  object_class->finalize = gst_hip_event_pool_finalize;
+}
+
+static void
+gst_hip_event_pool_init (GstHipEventPool * self)
+{
+  self->priv = new GstHipEventPoolPrivate ();
+}
+
+static void
+gst_hip_event_pool_finalize (GObject * object)
+{
+  auto self = GST_HIP_EVENT_POOL (object);
+
+  delete self->priv;
+
+  G_OBJECT_CLASS (parent_class)->finalize (object);
+}
+
+GstHipEventPool *
+gst_hip_event_pool_new (GstHipVendor vendor, guint device_id)
+{
+  g_return_val_if_fail (vendor != GST_HIP_VENDOR_UNKNOWN, nullptr);
+
+  auto self = (GstHipEventPool *)
+      g_object_new (GST_TYPE_HIP_EVENT_POOL, nullptr);
+  gst_object_ref_sink (self);
+
+  auto priv = self->priv;
+  priv->vendor = vendor;
+  priv->device_id = device_id;
+
+  return self;
+}
+
+static void
+gst_hip_event_pool_release (GstHipEventPool * pool, GstHipEvent * event)
+{
+  auto priv = pool->priv;
+  {
+    std::lock_guard < std::mutex > lk (priv->lock);
+    event->dispose = nullptr;
+    event->pool = nullptr;
+    priv->event_pool.push (event);
+  }
+
+  gst_object_unref (pool);
+}
+
+static gboolean
+gst_hip_event_dispose (GstHipEvent * event)
+{
+  if (!event->pool)
+    return TRUE;
+
+  gst_mini_object_ref (event);
+  gst_hip_event_pool_release (event->pool, event);
+
+  return FALSE;
+}
+
+static void
+gst_hip_event_free (GstHipEvent * event)
+{
+  delete event;
+}
+
+gboolean
+gst_hip_event_pool_acquire (GstHipEventPool * pool, GstHipEvent ** event)
+{
+  g_return_val_if_fail (GST_IS_HIP_EVENT_POOL (pool), FALSE);
+  g_return_val_if_fail (event, FALSE);
+
+  *event = nullptr;
+
+  auto priv = pool->priv;
+  GstHipEvent *new_event = nullptr;
+
+  {
+    std::lock_guard < std::mutex > lk (priv->lock);
+    if (!priv->event_pool.empty ()) {
+      new_event = priv->event_pool.front ();
+      priv->event_pool.pop ();
+    }
+  }
+
+  if (!new_event) {
+    auto hip_ret = HipSetDevice (priv->vendor, priv->device_id);
+    if (!gst_hip_result (hip_ret, priv->vendor)) {
+      GST_ERROR_OBJECT (pool, "Couldn't set device");
+      return FALSE;
+    }
+
+    hipEvent_t handle;
+    hip_ret = HipEventCreateWithFlags (priv->vendor, &handle,
+        hipEventDisableTiming);
+
+    if (!gst_hip_result (hip_ret, priv->vendor)) {
+      GST_ERROR_OBJECT (pool, "Couldn't create event");
+      return FALSE;
+    }
+
+    new_event = new GstHipEvent ();
+    new_event->handle = handle;
+    new_event->vendor = priv->vendor;
+    new_event->device_id = priv->device_id;
+
+    gst_mini_object_init (new_event, 0, gst_hip_event_get_type (),
+        nullptr, nullptr, (GstMiniObjectFreeFunction) gst_hip_event_free);
+  }
+
+  new_event->pool = (GstHipEventPool *) gst_object_ref (pool);
+  new_event->dispose = (GstMiniObjectDisposeFunction) gst_hip_event_dispose;
+
+  *event = new_event;
+
+  return TRUE;
+}
+
+GstHipVendor
+gst_hip_event_get_vendor (GstHipEvent * event)
+{
+  g_return_val_if_fail (event, GST_HIP_VENDOR_UNKNOWN);
+
+  return event->vendor;
+}
+
+guint
+gst_hip_event_get_device_id (GstHipEvent * event)
+{
+  g_return_val_if_fail (event, G_MAXUINT);
+
+  return event->vendor;
+}
+
+hipError_t
+gst_hip_event_record (GstHipEvent * event, hipStream_t stream)
+{
+  g_return_val_if_fail (event, hipErrorInvalidValue);
+
+  auto hip_ret = HipSetDevice (event->vendor, event->device_id);
+  if (!gst_hip_result (hip_ret, event->vendor))
+    return hip_ret;
+
+  return HipEventRecord (event->vendor, event->handle, stream);
+}
+
+hipError_t
+gst_hip_event_query (GstHipEvent * event)
+{
+  g_return_val_if_fail (event, hipErrorInvalidValue);
+
+  auto hip_ret = HipSetDevice (event->vendor, event->device_id);
+  if (!gst_hip_result (hip_ret, event->vendor))
+    return hip_ret;
+
+  return HipEventQuery (event->vendor, event->handle);
+}
+
+hipError_t
+gst_hip_event_synchronize (GstHipEvent * event)
+{
+  g_return_val_if_fail (event, hipErrorInvalidValue);
+
+  auto hip_ret = HipSetDevice (event->vendor, event->device_id);
+  if (!gst_hip_result (hip_ret, event->vendor))
+    return hip_ret;
+
+  return HipEventSynchronize (event->vendor, event->handle);
+}
+
+GstHipEvent *
+gst_hip_event_ref (GstHipEvent * event)
+{
+  return (GstHipEvent *) gst_mini_object_ref (event);
+}
+
+void
+gst_hip_event_unref (GstHipEvent * event)
+{
+  return gst_mini_object_unref (event);
+}
+
+void
+gst_clear_hip_event (GstHipEvent ** event)
+{
+  gst_clear_mini_object (event);
+}

subprojects/gst-plugins-bad/sys/hip/gsthipevent.h

@@ -0,0 +1,82 @@
+/* GStreamer
+ * Copyright (C) 2025 Seungha Yang <seungha@centricular.com>
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Library General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Library General Public License for more details.
+ *
+ * You should have received a copy of the GNU Library General Public
+ * License along with this library; if not, write to the
+ * Free Software Foundation, Inc., 51 Franklin St, Fifth Floor,
+ * Boston, MA 02110-1301, USA.
+ */
+
+#pragma once
+
+#include <gst/gst.h>
+#include "gsthip_fwd.h"
+#include "gsthip-enums.h"
+#include <hip/hip_runtime.h>
+
+G_BEGIN_DECLS
+
+#define GST_TYPE_HIP_EVENT_POOL                (gst_hip_event_pool_get_type ())
+#define GST_HIP_EVENT_POOL(obj)                (G_TYPE_CHECK_INSTANCE_CAST ((obj), GST_TYPE_HIP_EVENT_POOL, GstHipEventPool))
+#define GST_HIP_EVENT_POOL_CLASS(klass)        (G_TYPE_CHECK_CLASS_CAST ((klass), GST_TYPE_HIP_EVENT_POOL, GstHipEventPoolClass))
+#define GST_IS_HIP_EVENT_POOL(obj)             (G_TYPE_CHECK_INSTANCE_TYPE ((obj), GST_TYPE_HIP_EVENT_POOL))
+#define GST_IS_HIP_EVENT_POOL_CLASS(klass)     (G_TYPE_CHECK_CLASS_TYPE ((klass), GST_TYPE_HIP_EVENT_POOL))
+#define GST_HIP_EVENT_POOL_GET_CLASS(obj)      (G_TYPE_INSTANCE_GET_CLASS ((obj), GST_TYPE_HIP_EVENT_POOL, GstHipEventPoolClass))
+#define GST_HIP_EVENT_POOL_CAST(obj)           ((GstHipEventPool*)(obj))
+
+struct _GstHipEventPool
+{
+  GstObject parent;
+
+  /*< private >*/
+  GstHipEventPoolPrivate *priv;
+  gpointer _gst_reserved[GST_PADDING];
+};
+
+struct _GstHipEventPoolClass
+{
+  GstObjectClass parent_class;
+
+  /*< private >*/
+  gpointer _gst_reserved[GST_PADDING];
+};
+
+GType gst_hip_event_pool_get_type (void);
+
+GType gst_hip_event_get_type (void);
+
+GstHipEventPool * gst_hip_event_pool_new (GstHipVendor vendor,
+                                          guint device_id);
+
+gboolean          gst_hip_event_pool_acquire (GstHipEventPool * pool,
+                                              GstHipEvent ** event);
+
+GstHipVendor      gst_hip_event_get_vendor (GstHipEvent * event);
+
+guint             gst_hip_event_get_device_id (GstHipEvent * event);
+
+hipError_t        gst_hip_event_record (GstHipEvent * event,
+                                        hipStream_t stream);
+
+hipError_t        gst_hip_event_query (GstHipEvent * event);
+
+hipError_t        gst_hip_event_synchronize (GstHipEvent * event);
+
+GstHipEvent *     gst_hip_event_ref (GstHipEvent * event);
+
+void              gst_hip_event_unref (GstHipEvent * event);
+
+void              gst_clear_hip_event (GstHipEvent ** event);
+
+G_END_DECLS
+

subprojects/gst-plugins-bad/sys/hip/gsthiploader.cpp

@@ -73,6 +73,11 @@ struct GstHipFuncTableAmd
   hipError_t (*hipStreamCreate) (hipStream_t* stream);
   hipError_t (*hipStreamDestroy) (hipStream_t stream);
   hipError_t (*hipStreamSynchronize) (hipStream_t stream);
+  hipError_t (*hipEventCreateWithFlags) (hipEvent_t* event, unsigned flags);
+  hipError_t (*hipEventRecord) (hipEvent_t event, hipStream_t stream);
+  hipError_t (*hipEventDestroy) (hipEvent_t event);
+  hipError_t (*hipEventSynchronize) (hipEvent_t event);
+  hipError_t (*hipEventQuery) (hipEvent_t event);
   hipError_t (*hipModuleLoadData) (hipModule_t * module, const void *image);
   hipError_t (*hipModuleUnload) (hipModule_t module);
   hipError_t (*hipModuleGetFunction) (hipFunction_t * function,
@@ -163,6 +168,12 @@ struct GstHipFuncTableCudaRt
   cudaError_t (CUDAAPI *cudaStreamCreate) (cudaStream_t *pStream);
   cudaError_t (CUDAAPI *cudaStreamDestroy) (cudaStream_t stream);
   cudaError_t (CUDAAPI *cudaStreamSynchronize) (cudaStream_t stream);
+  cudaError_t (CUDAAPI *cudaEventCreateWithFlags) (cudaEvent_t *event,
+    unsigned int flags);
+  cudaError_t (CUDAAPI *cudaEventRecord) (cudaEvent_t event, cudaStream_t stream);
+  cudaError_t (CUDAAPI *cudaEventDestroy) (cudaEvent_t event);
+  cudaError_t (CUDAAPI *cudaEventSynchronize)(cudaEvent_t event);
+  cudaError_t (CUDAAPI *cudaEventQuery) (cudaEvent_t event);
   cudaError_t (CUDAAPI *cudaGraphicsMapResources) (int count,
     cudaGraphicsResource_t *resources, cudaStream_t stream);
   cudaError_t (CUDAAPI *cudaGraphicsResourceGetMappedPointer) (void **devPtr,
@@ -249,6 +260,11 @@ load_amd_func_table (void)
   LOAD_SYMBOL (hipStreamCreate);
   LOAD_SYMBOL (hipStreamDestroy);
   LOAD_SYMBOL (hipStreamSynchronize);
+  LOAD_SYMBOL (hipEventCreateWithFlags);
+  LOAD_SYMBOL (hipEventRecord);
+  LOAD_SYMBOL (hipEventDestroy);
+  LOAD_SYMBOL (hipEventSynchronize);
+  LOAD_SYMBOL (hipEventQuery);
   LOAD_SYMBOL (hipModuleLoadData);
   LOAD_SYMBOL (hipModuleUnload);
   LOAD_SYMBOL (hipModuleGetFunction);
@@ -370,6 +386,11 @@ load_cudart_func_table (guint major_ver, guint minor_ver)
   LOAD_SYMBOL (cudaStreamCreate);
   LOAD_SYMBOL (cudaStreamDestroy);
   LOAD_SYMBOL (cudaStreamSynchronize);
+  LOAD_SYMBOL (cudaEventCreateWithFlags);
+  LOAD_SYMBOL (cudaEventRecord);
+  LOAD_SYMBOL (cudaEventDestroy);
+  LOAD_SYMBOL (cudaEventSynchronize);
+  LOAD_SYMBOL (cudaEventQuery);
   LOAD_SYMBOL (cudaGraphicsMapResources);
   LOAD_SYMBOL (cudaGraphicsResourceGetMappedPointer);
   LOAD_SYMBOL (cudaGraphicsUnmapResources);
@@ -976,6 +997,68 @@ HipStreamSynchronize (GstHipVendor vendor, hipStream_t stream)
   return hipCUDAErrorTohipError (cuda_ret);
 }
 
+hipError_t
+HipEventCreateWithFlags (GstHipVendor vendor, hipEvent_t * event,
+    unsigned flags)
+{
+  CHECK_VENDOR (vendor);
+
+  if (vendor == GST_HIP_VENDOR_AMD)
+    return amd_ftable.hipEventCreateWithFlags (event, flags);
+
+  auto cuda_ret = cudart_ftable.cudaEventCreateWithFlags ((cudaEvent_t *) event,
+      flags);
+  return hipCUDAErrorTohipError (cuda_ret);
+}
+
+hipError_t
+HipEventRecord (GstHipVendor vendor, hipEvent_t event, hipStream_t stream)
+{
+  CHECK_VENDOR (vendor);
+
+  if (vendor == GST_HIP_VENDOR_AMD)
+    return amd_ftable.hipEventRecord (event, stream);
+
+  auto cuda_ret = cudart_ftable.cudaEventRecord ((cudaEvent_t) event, stream);
+  return hipCUDAErrorTohipError (cuda_ret);
+}
+
+hipError_t
+HipEventDestroy (GstHipVendor vendor, hipEvent_t event)
+{
+  CHECK_VENDOR (vendor);
+
+  if (vendor == GST_HIP_VENDOR_AMD)
+    return amd_ftable.hipEventDestroy (event);
+
+  auto cuda_ret = cudart_ftable.cudaEventDestroy ((cudaEvent_t) event);
+  return hipCUDAErrorTohipError (cuda_ret);
+}
+
+hipError_t
+HipEventSynchronize (GstHipVendor vendor, hipEvent_t event)
+{
+  CHECK_VENDOR (vendor);
+
+  if (vendor == GST_HIP_VENDOR_AMD)
+    return amd_ftable.hipEventSynchronize (event);
+
+  auto cuda_ret = cudart_ftable.cudaEventSynchronize ((cudaEvent_t) event);
+  return hipCUDAErrorTohipError (cuda_ret);
+}
+
+hipError_t
+HipEventQuery (GstHipVendor vendor, hipEvent_t event)
+{
+  CHECK_VENDOR (vendor);
+
+  if (vendor == GST_HIP_VENDOR_AMD)
+    return amd_ftable.hipEventQuery (event);
+
+  auto cuda_ret = cudart_ftable.cudaEventQuery ((cudaEvent_t) event);
+  return hipCUDAErrorTohipError (cuda_ret);
+}
+
 hipError_t
 HipModuleLoadData (GstHipVendor vendor, hipModule_t * module, const void *image)
 {

subprojects/gst-plugins-bad/sys/hip/gsthiploader.h

@@ -81,6 +81,23 @@ hipError_t HipStreamDestroy (GstHipVendor vendor,
 hipError_t HipStreamSynchronize (GstHipVendor vendor,
                                  hipStream_t stream);
 
+hipError_t HipEventCreateWithFlags (GstHipVendor vendor,
+                                    hipEvent_t* event,
+                                    unsigned flags);
+
+hipError_t HipEventRecord (GstHipVendor vendor,
+                           hipEvent_t event,
+                           hipStream_t stream);
+
+hipError_t HipEventDestroy (GstHipVendor vendor,
+                            hipEvent_t event);
+
+hipError_t HipEventSynchronize (GstHipVendor vendor,
+                                hipEvent_t event);
+
+hipError_t HipEventQuery (GstHipVendor vendor,
+                          hipEvent_t event);
+
 hipError_t HipModuleLoadData (GstHipVendor vendor,
                               hipModule_t* module,
                               const void* image);

subprojects/gst-plugins-bad/sys/hip/gsthipstream.cpp

@@ -51,9 +51,12 @@ struct _GstHipStream : public GstMiniObject
       if (gst_hip_result (hip_ret, vendor))
         HipStreamDestroy (vendor, handle);
     }
+
+    gst_clear_object (&event_pool);
   }
 
   hipStream_t handle = nullptr;
+  GstHipEventPool *event_pool = nullptr;
   GstHipVendor vendor;
   guint device_id;
 };
@@ -89,6 +92,7 @@ gst_hip_stream_new (GstHipVendor vendor, guint device_id)
   stream->handle = handle;
   stream->vendor = vendor;
   stream->device_id = device_id;
+  stream->event_pool = gst_hip_event_pool_new (vendor, device_id);
 
   gst_mini_object_init (stream, 0, gst_hip_stream_get_type (),
       nullptr, nullptr, (GstMiniObjectFreeFunction) gst_hip_stream_free);
@@ -121,6 +125,36 @@ gst_hip_stream_get_handle (GstHipStream * stream)
   return stream->handle;
 }
 
+gboolean
+gst_hip_stream_record_event (GstHipStream * stream, GstHipEvent ** event)
+{
+  g_return_val_if_fail (stream, FALSE);
+  g_return_val_if_fail (event, FALSE);
+
+  auto hip_ret = HipSetDevice (stream->vendor, stream->device_id);
+  if (!gst_hip_result (hip_ret, stream->vendor)) {
+    GST_ERROR ("Couldn't set device");
+    return FALSE;
+  }
+
+  GstHipEvent *new_event;
+  if (!gst_hip_event_pool_acquire (stream->event_pool, &new_event)) {
+    GST_ERROR ("Couldn't acquire event");
+    return FALSE;
+  }
+
+  hip_ret = gst_hip_event_record (new_event, stream->handle);
+  if (!gst_hip_result (hip_ret, stream->vendor)) {
+    GST_ERROR ("Couldn't record event");
+    gst_hip_event_unref (new_event);
+    return FALSE;
+  }
+
+  *event = new_event;
+
+  return TRUE;
+}
+
 GstHipStream *
 gst_hip_stream_ref (GstHipStream * stream)
 {

subprojects/gst-plugins-bad/sys/hip/gsthipstream.h

@@ -37,6 +37,9 @@ guint          gst_hip_stream_get_device_id (GstHipStream * stream);
 
 hipStream_t    gst_hip_stream_get_handle (GstHipStream * stream);
 
+gboolean       gst_hip_stream_record_event (GstHipStream * stream,
+                                            GstHipEvent ** event);
+
 GstHipStream * gst_hip_stream_ref (GstHipStream * stream);
 
 void           gst_hip_stream_unref (GstHipStream * stream);

subprojects/gst-plugins-bad/sys/hip/meson.build

@@ -13,6 +13,7 @@ hip_sources = [
   'gsthip-interop.cpp',
   'gsthipcompositor.cpp',
   'gsthipstream.cpp',
+  'gsthipevent.cpp',
   'plugin.cpp',
 ]
 

subprojects/gst-plugins-bad/sys/hip/stub/driver_types.h

@@ -399,6 +399,7 @@ typedef gpointer cudaStream_t;
 struct cudaGraphicsResource;
 
 typedef struct cudaGraphicsResource *cudaGraphicsResource_t;
+typedef struct CUevent_st *cudaEvent_t;
 
 G_END_DECLS
 

subprojects/gst-plugins-bad/sys/hip/stub/hip/hip_runtime_api.h

@@ -483,4 +483,24 @@ typedef hipGraphicsResource* hipGraphicsResource_t;
 typedef struct ihipStream_t* hipStream_t;
 typedef struct ihipModule_t* hipModule_t;
 typedef struct ihipModuleSymbol_t* hipFunction_t;
+typedef struct ihipEvent_t* hipEvent_t;
+
+/** Default stream creation flags. These are used with hipStreamCreate().*/
+#define hipStreamDefault  0x00
+
+/** Stream does not implicitly synchronize with null stream.*/
+#define hipStreamNonBlocking 0x01
+
+//Flags that can be used with hipEventCreateWithFlags.
+/** Default flags.*/
+#define hipEventDefault 0x0
+
+/** Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency.*/
+#define hipEventBlockingSync 0x1
+
+/** Disable event's capability to record timing information. May improve performance.*/
+#define hipEventDisableTiming  0x2
+
+/** Event can support IPC. hipEventDisableTiming also must be set.*/
+#define hipEventInterprocess 0x4