Bug fix: events would get created on the wrong device in a multi-gpu scenario. (#3910)

Currently the GPU events are being created without setting the device, which would thus fail for a multi-gpu setup.

Author: lakhinderwalia

src/targets/gpu/hip.cpp

@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -237,7 +237,7 @@ void set_device(std::size_t id)
 {
     auto status = hipSetDevice(id);
     if(status != hipSuccess)
-        MIGRAPHX_THROW("Error setting device");
+        MIGRAPHX_THROW("Error setting device: " + hip_error(status));
 }
 
 void gpu_sync()

src/targets/gpu/include/migraphx/gpu/context.hpp

@@ -59,7 +59,10 @@ struct hip_device
     {
         auto status = hipGetDeviceProperties(&device_props, device_id);
         if(status != hipSuccess)
-            MIGRAPHX_THROW("Failed to allocate stream");
+            MIGRAPHX_THROW("Failed to get device properties: " + hip_error(status));
+
+        // Set the device prior to Events that get created within a Context.
+        set_device(device_id);
 
         for(std::size_t i = 0; i < n; i++)
             add_stream();
@@ -80,7 +83,7 @@ struct hip_device
             hipStream_t result = nullptr;
             auto status        = hipStreamCreateWithFlags(&result, hipStreamNonBlocking);
             if(status != hipSuccess)
-                MIGRAPHX_THROW("Failed to allocate stream");
+                MIGRAPHX_THROW("Failed to allocate stream: " + hip_error(status));
             return hip_stream_ptr{result};
         }
 
@@ -147,23 +150,23 @@ struct hip_device
             setup();
             auto status = hipStreamSynchronize(s.get());
             if(status != hipSuccess)
-                MIGRAPHX_THROW("Failed to wait.");
+                MIGRAPHX_THROW("Failed to wait: " + hip_error(status));
         }
 
         void wait(hipEvent_t event)
         {
             setup();
             auto status = hipStreamWaitEvent(get(), event, 0);
             if(status != hipSuccess)
-                MIGRAPHX_THROW("Failed to wait.");
+                MIGRAPHX_THROW("Failed to wait: " + hip_error(status));
         }
 
         void record(hipEvent_t event)
         {
             setup();
             auto status = hipEventRecord(event, get());
             if(status != hipSuccess)
-                MIGRAPHX_THROW("Failed to record.");
+                MIGRAPHX_THROW("Failed to record: " + hip_error(status));
         }
 
         private:
@@ -293,7 +296,7 @@ struct context
         hipEvent_t event;
         auto status = hipEventCreateWithFlags(&event, hipEventDisableTiming);
         if(status != hipSuccess)
-            MIGRAPHX_THROW("Failed to create event");
+            MIGRAPHX_THROW("Failed to create event: " + hip_error(status));
         return hip_event_ptr{event};
     }
 
@@ -302,7 +305,7 @@ struct context
         hipEvent_t event;
         auto status = hipEventCreate(&event);
         if(status != hipSuccess)
-            MIGRAPHX_THROW("Failed to create event");
+            MIGRAPHX_THROW("Failed to create event: " + hip_error(status));
         return hip_event_ptr{event};
     }
 
@@ -332,7 +335,7 @@ struct context
     {
         auto status = hipEventRecord(begin_event.get(), queue.get<hipStream_t>());
         if(status != hipSuccess)
-            MIGRAPHX_THROW("failed to record " + hip_error(status));
+            MIGRAPHX_THROW("Failed to record: " + hip_error(status));
 
         get_stream().wait(begin_event.get());
     }
@@ -343,7 +346,7 @@ struct context
 
         auto status = hipStreamWaitEvent(queue.get<hipStream_t>(), finish_event.get(), 0);
         if(status != hipSuccess)
-            MIGRAPHX_THROW("Failed to wait on event " + hip_error(status));
+            MIGRAPHX_THROW("Failed to wait on event: " + hip_error(status));
     }
 
     any_ptr get_queue() { return get_stream().get(); }
@@ -383,9 +386,9 @@ struct context
     // for event perf timing
     shared<hip_event_ptr> start_event = nullptr;
     shared<hip_event_ptr> stop_event  = nullptr;
-    // for stream syncronization
-    shared<hip_event_ptr> begin_event  = nullptr;
-    shared<hip_event_ptr> finish_event = nullptr;
+    // for stream synchronization
+    shared<hip_event_ptr> begin_event           = nullptr;
+    shared<hip_event_ptr> finish_event          = nullptr;
     std::shared_ptr<auto_save_problem_cache> pc = nullptr;
 };