From 3e6d1858af55c68a51af46ced97d4f1d0b0a35f9 Mon Sep 17 00:00:00 2001
From: shawnz <shawnz@nvidia.com>
Date: Fri, 5 Sep 2025 11:55:32 +0800
Subject: [PATCH] Bug 5502929: Update systemWideAtomics to add device sync for
 WSL

---
 .../systemWideAtomics/systemWideAtomics.cu               | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/Samples/0_Introduction/systemWideAtomics/systemWideAtomics.cu b/Samples/0_Introduction/systemWideAtomics/systemWideAtomics.cu
index 121ba207..4b7e46b1 100644
--- a/Samples/0_Introduction/systemWideAtomics/systemWideAtomics.cu
+++ b/Samples/0_Introduction/systemWideAtomics/systemWideAtomics.cu
@@ -297,6 +297,9 @@ int main(int argc, char **argv)
         exit(EXIT_WAIVED);
     }
 
+    int concurrentManagedAccess;
+    checkCudaErrors(cudaDeviceGetAttribute(&concurrentManagedAccess, cudaDevAttrConcurrentManagedAccess, dev_id));
+
     if (device_prop.major < 6) {
         printf("%s: requires a minimum CUDA compute 6.0 capability, waiving "
                "testing.\n",
@@ -326,6 +329,12 @@ int main(int argc, char **argv)
     atom_arr[7] = atom_arr[9] = 0xff;
 
     atomicKernel<<<numBlocks, numThreads>>>(atom_arr);
+
+    // Do device synchronize for devices that don't support concurrent managed access such as WSL.
+    if (!concurrentManagedAccess) {
+        checkCudaErrors(cudaDeviceSynchronize());
+    }
+
     atomicKernel_CPU(atom_arr, numBlocks * numThreads);
 
     checkCudaErrors(cudaDeviceSynchronize());