diff --git a/Common/helper_multiprocess.cpp b/Common/helper_multiprocess.cpp
index c7769b50..ee1618aa 100644
--- a/Common/helper_multiprocess.cpp
+++ b/Common/helper_multiprocess.cpp
@@ -188,7 +188,7 @@ int ipcCreateSocket(ipcHandle *&handle, const char *name,
   char path_name[50];
 
   // Create unique name for the socket with path if SOCK_FOLDER is set.
-  sprintf(path_name, "%s%u", SOCK_FOLDER, getpid());
+  sprintf(path_name, "%s/%u", getSocketFolder().c_str(), getpid());
 
   unlink(path_name);
   memset(&servaddr, 0, sizeof(servaddr));
@@ -230,7 +230,7 @@ int ipcOpenSocket(ipcHandle *&handle) {
   char temp[50];
 
   // Create unique name for the socket with path if SOCK_FOLDER is set.
-  sprintf(temp, "%s%u", SOCK_FOLDER, getpid());
+  sprintf(temp, "%s/%u", getSocketFolder().c_str(), getpid());
 
   strcpy(cliaddr.sun_path, temp);
   if (bind(sock, (struct sockaddr *)&cliaddr, sizeof(cliaddr)) < 0) {
@@ -367,7 +367,7 @@ int ipcSendShareableHandle(ipcHandle *handle,
   memset(&cliaddr, 0, sizeof(cliaddr));
   cliaddr.sun_family = AF_UNIX;
   char temp[20];
-  sprintf(temp, "%s%u", SOCK_FOLDER, process);
+  sprintf(temp, "%s/%u", getSocketFolder().c_str(), process);
   strcpy(cliaddr.sun_path, temp);
   len = sizeof(cliaddr);
 
diff --git a/Common/helper_multiprocess.h b/Common/helper_multiprocess.h
index d0bdcc20..2c09b5e3 100644
--- a/Common/helper_multiprocess.h
+++ b/Common/helper_multiprocess.h
@@ -52,14 +52,12 @@
 #include <memory.h>
 #include <sys/un.h>
 #endif
+#include <filesystem>
 #include <vector>
 
-// Define "/tmp" as socket creating folder for QNX
-#if defined(__QNX__)
-#define SOCK_FOLDER "/tmp/" 
-#else
-#define SOCK_FOLDER ""
-#endif
+inline std::string getSocketFolder() {
+    return std::filesystem::temp_directory_path().string();
+}
 
 typedef struct sharedMemoryInfo_st {
     void *addr;
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());