NASP/cuda-samples
53
0
Fork 0
mirror of https://github.com/NVIDIA/cuda-samples.git synced 2026-01-09 11:22:51 +08:00
Code Issues Packages Projects Releases Wiki Activity

Bug 5502929: Update systemWideAtomics to add device sync for WSL

Browse Source
This commit is contained in:
shawnz 2025-09-05 11:55:32 +08:00
parent 12b046a80e
commit 3e6d1858af
1 changed files with 9 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
Samples/0_Introduction/systemWideAtomics/systemWideAtomics.cu
View File

@ -297,6 +297,9 @@ int main(int argc, char **argv)
exit(EXIT_WAIVED); exit(EXIT_WAIVED);
} }
int concurrentManagedAccess;
checkCudaErrors(cudaDeviceGetAttribute(&concurrentManagedAccess, cudaDevAttrConcurrentManagedAccess, dev_id));
if (device_prop.major < 6) { if (device_prop.major < 6) {
printf("%s: requires a minimum CUDA compute 6.0 capability, waiving " printf("%s: requires a minimum CUDA compute 6.0 capability, waiving "
"testing.\n", "testing.\n",
@ -326,6 +329,12 @@ int main(int argc, char **argv)
atom_arr[7] = atom_arr[9] = 0xff; atom_arr[7] = atom_arr[9] = 0xff;
atomicKernel<<<numBlocks, numThreads>>>(atom_arr); 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); atomicKernel_CPU(atom_arr, numBlocks * numThreads);
checkCudaErrors(cudaDeviceSynchronize()); checkCudaErrors(cudaDeviceSynchronize());