Fix cross device synchronization (#2693)

Co-authored-by: mayuyuace <[email protected]>

Author: LuFinch

WORKSPACE

@@ -32,7 +32,7 @@ load(
 #    curl -L https://github.com/openxla/xla/archive/<git hash>.tar.gz | sha256sum
 #    and update XLA_SHA256 with the result.
 
-XLA_COMMIT = "98bf412acc77c515ed7b40cff26ac1f1decf8f5b"
+XLA_COMMIT = "da290b48daafac8cf9db1d9431ed1e4bd5fccdfa"
 
 new_git_repository(
     name = "intel_extension_for_openxla",

itex/core/devices/gpu/gpu_device_plugin.cc

@@ -166,32 +166,28 @@ void gpu_destroy_event(const SP_Device* device, SP_Event event) {
 
 // Requests the current status of the event from the underlying platform.
 SE_EventStatus gpu_get_event_status(const SP_Device* device, SP_Event event) {
-  if (IsMultipleStreamEnabled()) {
-    ITEX_GPUEvent event_handle = static_cast<SP_Event_st*>(event)->event_handle;
-    auto event_status =
-        event_handle.get_info<sycl::info::event::command_execution_status>();
-    switch (event_status) {
-      case sycl::info::event_command_status::submitted:
-      case sycl::info::event_command_status::running:
-        return SE_EVENT_PENDING;
-      case sycl::info::event_command_status::complete:
-        return SE_EVENT_COMPLETE;
-      default:
-        return SE_EVENT_UNKNOWN;
-    }
+  ITEX_GPUEvent event_handle = static_cast<SP_Event_st*>(event)->event_handle;
+  auto event_status =
+      event_handle.get_info<sycl::info::event::command_execution_status>();
+  switch (event_status) {
+    case sycl::info::event_command_status::submitted:
+    case sycl::info::event_command_status::running:
+      return SE_EVENT_PENDING;
+    case sycl::info::event_command_status::complete:
+      return SE_EVENT_COMPLETE;
+    default:
+      return SE_EVENT_UNKNOWN;
   }
   return SE_EVENT_COMPLETE;
 }
 
 // Inserts the specified event at the end of the specified stream.
 void gpu_record_event(const SP_Device* device, SP_Stream stream, SP_Event event,
                       TF_Status* status) {
-  if (IsMultipleStreamEnabled()) {
-    ITEX_GPUStream* stream_handle =
-        static_cast<SP_Stream_st*>(stream)->stream_handle;
-    ITEX_GPUEvent recorded_event = stream_handle->ext_oneapi_submit_barrier();
-    event->event_handle = recorded_event;
-  }
+  ITEX_GPUStream* stream_handle =
+      static_cast<SP_Stream_st*>(stream)->stream_handle;
+  ITEX_GPUEvent recorded_event = stream_handle->ext_oneapi_submit_barrier();
+  event->event_handle = recorded_event;
 }
 
 // Wait for the specified event at the end of the specified stream.

itex/core/devices/gpu/itex_gpu_runtime.cc

@@ -365,25 +365,16 @@ ITEX_GPUError_t ITEX_GPUDestroyEvent(ITEX_GPUDevice* device_handle,
 
 ITEX_GPUError_t ITEX_GPUStreamWaitEvent(ITEX_GPUStream* stream,
                                         ITEX_GPUEvent event) {
-  if (IsMultipleStreamEnabled()) {
-    const std::vector<ITEX_GPUEvent> event_list{event};
-    stream->ext_oneapi_submit_barrier(event_list);
-  } else {
-    stream->wait();
-  }
+  const std::vector<ITEX_GPUEvent> event_list{event};
+  stream->ext_oneapi_submit_barrier(event_list);
   return ITEX_GPU_SUCCESS;
 }
 
 ITEX_GPUError_t ITEX_GPUStreamWaitStream(ITEX_GPUStream* dependent,
                                          ITEX_GPUStream* other) {
-  if (IsMultipleStreamEnabled()) {
-    ITEX_GPUEvent event = other->ext_oneapi_submit_barrier();
-    const std::vector<ITEX_GPUEvent> event_list{event};
-    dependent->ext_oneapi_submit_barrier(event_list);
-  } else {
-    dependent->wait();
-    other->wait();
-  }
+  ITEX_GPUEvent event = other->ext_oneapi_submit_barrier();
+  const std::vector<ITEX_GPUEvent> event_list{event};
+  dependent->ext_oneapi_submit_barrier(event_list);
   return ITEX_GPU_SUCCESS;
 }