diff --git a/src/gallium/frontends/rusticl/api/icd.rs b/src/gallium/frontends/rusticl/api/icd.rs
index 4844c238956..b0838fbb971 100644
--- a/src/gallium/frontends/rusticl/api/icd.rs
+++ b/src/gallium/frontends/rusticl/api/icd.rs
@@ -84,9 +84,9 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
     clEnqueueNDRangeKernel: Some(cl_enqueue_ndrange_kernel),
     clEnqueueTask: Some(cl_enqueue_task),
     clEnqueueNativeKernel: None,
-    clEnqueueMarker: None,
+    clEnqueueMarker: Some(cl_enqueue_marker),
     clEnqueueWaitForEvents: None,
-    clEnqueueBarrier: None,
+    clEnqueueBarrier: Some(cl_enqueue_barrier),
     clGetExtensionFunctionAddress: Some(cl_get_extension_function_address),
     clCreateFromGLBuffer: None,
     clCreateFromGLTexture2D: None,
@@ -127,8 +127,8 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
     clEnqueueFillBuffer: Some(cl_enqueue_fill_buffer),
     clEnqueueFillImage: Some(cl_enqueue_fill_image),
     clEnqueueMigrateMemObjects: None,
-    clEnqueueMarkerWithWaitList: None,
-    clEnqueueBarrierWithWaitList: None,
+    clEnqueueMarkerWithWaitList: Some(cl_enqueue_marker_with_wait_list),
+    clEnqueueBarrierWithWaitList: Some(cl_enqueue_barrier_with_wait_list),
     clGetExtensionFunctionAddressForPlatform: None,
     clCreateFromGLTexture: None,
     clGetDeviceIDsFromD3D11KHR: ptr::null_mut(),
@@ -1130,6 +1130,14 @@ extern "C" fn cl_enqueue_task(
     ))
 }
 
+extern "C" fn cl_enqueue_marker(command_queue: cl_command_queue, event: *mut cl_event) -> cl_int {
+    match_err!(enqueue_marker(command_queue, event))
+}
+
+extern "C" fn cl_enqueue_barrier(command_queue: cl_command_queue) -> cl_int {
+    match_err!(enqueue_barrier(command_queue))
+}
+
 extern "C" fn cl_get_extension_function_address(
     function_name: *const ::std::os::raw::c_char,
 ) -> *mut ::std::ffi::c_void {
@@ -1412,6 +1420,34 @@ extern "C" fn cl_enqueue_fill_image(
     CL_OUT_OF_HOST_MEMORY
 }
 
+extern "C" fn cl_enqueue_marker_with_wait_list(
+    command_queue: cl_command_queue,
+    num_events_in_wait_list: cl_uint,
+    event_wait_list: *const cl_event,
+    event: *mut cl_event,
+) -> cl_int {
+    match_err!(enqueue_marker_with_wait_list(
+        command_queue,
+        num_events_in_wait_list,
+        event_wait_list,
+        event
+    ))
+}
+
+extern "C" fn cl_enqueue_barrier_with_wait_list(
+    command_queue: cl_command_queue,
+    num_events_in_wait_list: cl_uint,
+    event_wait_list: *const cl_event,
+    event: *mut cl_event,
+) -> cl_int {
+    match_err!(enqueue_barrier_with_wait_list(
+        command_queue,
+        num_events_in_wait_list,
+        event_wait_list,
+        event
+    ))
+}
+
 extern "C" fn cl_create_command_queue_with_properties(
     context: cl_context,
     device: cl_device_id,
diff --git a/src/gallium/frontends/rusticl/api/queue.rs b/src/gallium/frontends/rusticl/api/queue.rs
index ee2f09c7aa6..9da0f6bbe36 100644
--- a/src/gallium/frontends/rusticl/api/queue.rs
+++ b/src/gallium/frontends/rusticl/api/queue.rs
@@ -1,7 +1,9 @@
 extern crate rusticl_opencl_gen;
 
+use crate::api::event::create_and_queue;
 use crate::api::icd::*;
 use crate::api::util::*;
+use crate::core::event::*;
 use crate::core::queue::*;
 
 use self::rusticl_opencl_gen::*;
@@ -70,6 +72,69 @@ pub fn create_command_queue(
     Ok(cl_command_queue::from_arc(Queue::new(c, d, properties)?))
 }
 
+pub fn enqueue_marker(command_queue: cl_command_queue, event: *mut cl_event) -> CLResult<()> {
+    let q = command_queue.get_arc()?;
+
+    // TODO marker makes sure previous commands did complete
+    create_and_queue(
+        q,
+        CL_COMMAND_MARKER,
+        Vec::new(),
+        event,
+        false,
+        Box::new(|_, _| Ok(())),
+    )
+}
+
+pub fn enqueue_marker_with_wait_list(
+    command_queue: cl_command_queue,
+    num_events_in_wait_list: cl_uint,
+    event_wait_list: *const cl_event,
+    event: *mut cl_event,
+) -> CLResult<()> {
+    let q = command_queue.get_arc()?;
+    let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+    // TODO marker makes sure previous commands did complete
+    create_and_queue(
+        q,
+        CL_COMMAND_MARKER,
+        evs,
+        event,
+        false,
+        Box::new(|_, _| Ok(())),
+    )
+}
+
+pub fn enqueue_barrier(command_queue: cl_command_queue) -> CLResult<()> {
+    let q = command_queue.get_arc()?;
+
+    // TODO barriers make sure previous commands did complete and other commands didn't start
+    let e = Event::new(&q, CL_COMMAND_BARRIER, Vec::new(), Box::new(|_, _| Ok(())));
+    q.queue(&e);
+    Ok(())
+}
+
+pub fn enqueue_barrier_with_wait_list(
+    command_queue: cl_command_queue,
+    num_events_in_wait_list: cl_uint,
+    event_wait_list: *const cl_event,
+    event: *mut cl_event,
+) -> CLResult<()> {
+    let q = command_queue.get_arc()?;
+    let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+    // TODO barriers make sure previous commands did complete and other commands didn't start
+    create_and_queue(
+        q,
+        CL_COMMAND_BARRIER,
+        evs,
+        event,
+        false,
+        Box::new(|_, _| Ok(())),
+    )
+}
+
 pub fn flush_queue(command_queue: cl_command_queue) -> CLResult<()> {
     // CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
     command_queue.get_ref()?.flush(false)