Skip hipStreamSynchronize on iGPU when stream is idle

On integrated GPUs with fine-grained coherent unified memory,
hipStreamSynchronize is unnecessary when the stream has no pending
work. Use hipStreamQuery (non-blocking) to check first, only sync
when the stream is actually busy.

Results: hipStreamSynchronize calls reduced from 5683 to 54 (-99%).

Author: Geramy

mlx/backend/rocm/allocator.cpp

@@ -550,7 +550,12 @@ void* Buffer::raw_ptr() {
   auto& cbuf = *static_cast<rocm::RocmBuffer*>(ptr_);
 
   if (cbuf.device == -1) {
-    (void)hipStreamSynchronize(nullptr);
+    // Unified memory on iGPU: fine-grained coherent memory means CPU sees
+    // GPU writes without explicit sync. Only sync if the stream has pending
+    // work (hipStreamQuery returns hipErrorNotReady when busy).
+    if (hipStreamQuery(nullptr) != hipSuccess) {
+      (void)hipStreamSynchronize(nullptr);
+    }
   } else {
     (void)hipDeviceSynchronize();
     rocm::allocator().move_to_unified_memory(cbuf);