[ROCm] QDP: order default-stream setup before forked-stream kernels

Restore the non-blocking forked stream on all platforms and fix the
actual root cause of the Windows gfx1201 stream-test regression: a
missing ordering between default-stream buffer setup and a kernel
launched on a non-blocking stream that consumes it.

Background: fork_default_stream creates a hipStreamNonBlocking stream to
match cudarc and preserve the dual-stream copy/compute overlap. The
encoders set up their input (htod) and output (alloc_zeros) buffers with
the blocking shim copies, which run on the NULL/default stream, then
launch the norm/encode kernels on the caller's forked stream.

CUDA's legacy default stream is synchronizing, so on NVIDIA that setup is
implicitly ordered before the forked-stream kernel reads the buffer. HIP's
default stream is NOT synchronizing relative to a non-blocking stream, so
the kernel raced the setup and read stale/zero data; the norm came back as
the zero-initialized value and the result was wrong. A previous change
worked around this by creating a blocking stream on Windows, which masked
every such site at once but sacrificed the pipeline overlap and was
attributed to a nonexistent "cache coherency gap" in the Windows runtime.
A minimal HIP reproducer shows the runtime is correct: a kernel on a
non-blocking stream that is properly ordered against the default stream
(stream sync, device sync, or an event wait) reads back correct data on
both the Adrenalin and ROCm runtimes; only the unordered case fails.

The fix synchronizes the default stream at the end of the blocking
alloc_zeros / htod copies (sync_default_stream), restoring the
CUDA-equivalent ordering on every platform without touching the async
pipeline (which uses async copies on explicit streams, not these blocking
paths). The forked stream is non-blocking again everywhere. The symmetric
readback hazard (a NULL-stream dtoh after a forked-stream kernel) is closed
where it was still missing a sync: the f64/f32 batch norm-validation copies
in amplitude.rs and the phase batch finiteness-probe copy in phase.rs now
synchronize the caller's stream before reading back, matching the idiom the
other encoders already use.

This is arch-unified: correct on wave32 (gfx1100/gfx1151/gfx1201) and
wave64 (gfx90a). On Linux the default stream is already synchronizing, so
the added sync is a harmless no-op-cost ordering point.

Authored with the assistance of Claude (Anthropic).

Test Plan

gfx1201 (RX 9070 XT, Windows 11, TheRock ROCm 7.14, HIP_VISIBLE_DEVICES=0):

```
export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx1201 HIP_VISIBLE_DEVICES=0
cargo build -p qdp-core -p qdp-kernels --no-default-features --features hip
cargo test  -p qdp-core -p qdp-kernels --no-default-features --features hip -- --test-threads=1
```

All suites pass, 0 failures. The three previously-regressed stream tests
pass with the non-blocking stream restored:
test_l2_norm_batch_kernel_stream,
test_encode_from_gpu_ptr_f32_with_stream_non_default_success,
test_encode_batch_from_gpu_ptr_f32_with_stream_success. qdp-core lib 77,
gpu_ptr_encoding 68 (all 10 _with_stream variants pass), amplitude 21,
angle 10, and all other GPU and non-GPU suites green.

Author: jeffdaily

qdp/qdp-core/src/gpu/encodings/amplitude.rs

@@ -414,6 +414,12 @@ impl QuantumEncoder for AmplitudeEncoder {
         };
         {
             crate::profile_scope!("GPU::NormValidation");
+            // The norm kernel ran on the caller's stream, but dtoh_sync_copy reads
+            // back on the default stream. Synchronize the caller's stream first so
+            // the result is visible: with a non-blocking stream (which does not
+            // implicitly order against the default stream) the readback would
+            // otherwise race and observe the zero-initialized buffer.
+            sync_cuda_stream(stream, "Norm stream synchronize failed (batch)")?;
             let host_inv_norms = device
                 .dtoh_sync_copy(&inv_norms_gpu)
                 .map_err(|e| MahoutError::Cuda(format!("Failed to copy norms to host: {:?}", e)))?;
@@ -646,6 +652,12 @@ impl QuantumEncoder for AmplitudeEncoder {
         };
         {
             crate::profile_scope!("GPU::NormValidation_f32");
+            // The norm kernel ran on the caller's stream, but dtoh_sync_copy reads
+            // back on the default stream. Synchronize the caller's stream first so
+            // the result is visible: with a non-blocking stream (which does not
+            // implicitly order against the default stream) the readback would
+            // otherwise race and observe the zero-initialized buffer.
+            sync_cuda_stream(stream, "Norm stream synchronize failed (batch f32)")?;
             let host_inv_norms = device
                 .dtoh_sync_copy(&inv_norms_gpu)
                 .map_err(|e| MahoutError::Cuda(format!("Failed to copy norms to host: {:?}", e)))?;

qdp/qdp-core/src/gpu/encodings/phase.rs

@@ -353,6 +353,15 @@ impl QuantumEncoder for PhaseEncoder {
 
         {
             crate::profile_scope!("GPU::PhaseFiniteValidationHostCopy");
+            // The norm probe ran on the caller's stream, but dtoh_sync_copy reads
+            // back on the default stream. Synchronize the caller's stream first so
+            // the result is visible: with a non-blocking stream (which does not
+            // implicitly order against the default stream) the readback would
+            // otherwise race and observe the zero-initialized buffer.
+            crate::gpu::cuda_sync::sync_cuda_stream(
+                stream,
+                "Phase validation norm stream synchronize failed (batch)",
+            )?;
             let host_norms = device
                 .dtoh_sync_copy(&phase_validation_buffer)
                 .map_err(|e| {

qdp/qdp-kernels/src/device.rs

@@ -59,8 +59,6 @@ mod hip {
     const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
     // hipStreamNonBlocking: the new stream does not implicitly synchronize with
     // the NULL/default stream, matching cudarc's fork_default_stream.
-    // Used on Linux only; Windows uses a blocking stream (see fork_default_stream).
-    #[cfg(target_os = "linux")]
     const HIP_STREAM_NON_BLOCKING: u32 = 1;
 
     unsafe extern "C" {
@@ -71,10 +69,7 @@ mod hip {
         fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
         fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32) -> hipError_t;
         fn hipDeviceSynchronize() -> hipError_t;
-        #[cfg(target_os = "linux")]
         fn hipStreamCreateWithFlags(stream: *mut *mut c_void, flags: u32) -> hipError_t;
-        #[cfg(not(target_os = "linux"))]
-        fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
         fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
         fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
     }
@@ -92,6 +87,24 @@ mod hip {
         }
     }
 
+    /// Synchronize the NULL/default stream so its prior work is ordered before any
+    /// other stream observes the affected buffers.
+    ///
+    /// The blocking shim copies (htod/alloc_zeros) issue hipMemcpy/hipMemset on the
+    /// default (NULL) stream. CUDA's legacy default stream is synchronizing, so on
+    /// NVIDIA that work is implicitly ordered before a kernel launched on a forked
+    /// non-blocking stream that reads the same buffer. HIP's default stream is NOT
+    /// synchronizing relative to a hipStreamNonBlocking stream, so without this an
+    /// encoder that sets up input/output on the default stream and then launches
+    /// the norm/encode kernels on the caller's forked stream would race the setup
+    /// (the kernel reads stale/zero data). A default-stream synchronize after the
+    /// blocking copy restores the CUDA-equivalent ordering while preserving the
+    /// dual-stream copy/compute overlap (which uses async copies on explicit
+    /// streams, not these blocking paths).
+    fn sync_default_stream() -> Result<(), DriverError> {
+        unsafe { check(hipStreamSynchronize(std::ptr::null_mut())) }
+    }
+
     /// Marker: type is safe to byte-copy to/from the device. Mirrors
     /// `cudarc::driver::DeviceRepr`.
     ///
@@ -305,6 +318,7 @@ mod hip {
             if bytes > 0 {
                 unsafe {
                     check(hipMemset(slice.raw_ptr(), 0, bytes))?;
+                    sync_default_stream()?;
                 }
             }
             Ok(slice)
@@ -354,6 +368,7 @@ mod hip {
                         bytes,
                         HIP_MEMCPY_HOST_TO_DEVICE,
                     ))?;
+                    sync_default_stream()?;
                 }
             }
             Ok(())
@@ -400,23 +415,18 @@ mod hip {
             self.bind()?;
             let mut stream: *mut c_void = std::ptr::null_mut();
             unsafe {
-                // On Linux, use a non-blocking stream (HIP_STREAM_NON_BLOCKING=1) to
-                // match cudarc: a blocking stream serializes H2D copies against the NULL
-                // stream and defeats the dual-stream copy/compute overlap the pipeline
-                // relies on.
-                //
-                // On Windows (TheRock ROCm 7.14), non-blocking stream writes are not
-                // visible via hipMemcpy from the host after hipStreamSynchronize, due to
-                // a cache coherency gap in the Windows HIP runtime. Use a blocking stream
-                // to restore correct D2H readback semantics. The pipeline overlap
-                // optimization is not available on Windows ROCm.
-                #[cfg(target_os = "linux")]
+                // Non-blocking stream (matches cudarc on both Linux and Windows): a
+                // default/blocking stream would serialize H2D copies against the NULL
+                // stream and defeat the dual-stream copy/compute overlap the pipeline
+                // relies on. Because a non-blocking stream does not implicitly order
+                // against the default stream, any host-visible readback (a default/NULL
+                // stream dtoh copy) of results produced on this stream MUST first
+                // synchronize this stream (wait_for / sync_cuda_stream); the encoders
+                // already do so before every readback.
                 check(hipStreamCreateWithFlags(
                     &mut stream,
                     HIP_STREAM_NON_BLOCKING,
                 ))?;
-                #[cfg(not(target_os = "linux"))]
-                check(hipStreamCreate(&mut stream))?;
             }
             Ok(CudaStream {
                 stream,