Gate the pinned-memory system-scope atomic test on host native atomics

[maleadt]

The "Compute sanitizer" CI job has been failing on main since #3268 (first failing build 7824), with 32 reports of

Invalid __global__ system-scoped atomic of size 4 bytes

from the system scope testset in test/core/device/intrinsics/atomics.jl. The device frame points at the definition of add_kernel, but the flagged access is the host-pinned smoke test that #3268 added at the end of the testset (a Int32[0] wrapped as a HostMemory CuArray).

Compute-sanitizer 2026.3 checks system-scoped atomics against the memory they target (its internal message reads "A system-atomic operation was performed on non-migratable memory"). Per the programming guide's atomicity rules, a system-scope atomic on mapped/pinned memory is only atomic at system scope when CU_DEVICE_ATTRIBUTE_HOST_NATIVE_ATOMIC_SUPPORTED is set, which is not the case on x86 PCIe hosts such as the CI agents. Device and managed memory are not flagged. There is no sanitizer option to disable the check.

This gates the pinned-memory smoke test on that attribute, so it only runs where the operation is actually meaningful (e.g., NVLink C2C systems).

Verified locally on an RTX 5080 with NVPTX_LLVM_Backend_jll 23 (which emits atom.acq_rel.sys.global.add, unlike backend 22 that drops the scope and thus never triggered this locally): the sanitized atomics test reproduces the CI failure before this change and passes with no sanitizer findings after it.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 92.56%. Comparing base (1d1de2a) to head (90c0348).

Additional details and impacted files

@@            Coverage Diff             @@
##             main    #3271      +/-   ##
==========================================
- Coverage   92.58%   92.56%   -0.02%     
==========================================
  Files         188      188              
  Lines       16471    16472       +1     
==========================================
- Hits        15249    15247       -2     
- Misses       1222     1225       +3     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

[github-actions]

CUDA.jl Benchmarks

Details

Benchmark suite	Current: 90c0348	Previous: 1d1de2a	Ratio
array/accumulate/Float32/1d	98109 ns	98395 ns	1.00
array/accumulate/Float32/dims=1	71774 ns	71732 ns	1.00
array/accumulate/Float32/dims=1L	1589136 ns	1589534 ns	1.00
array/accumulate/Float32/dims=2	138335 ns	137727 ns	1.00
array/accumulate/Float32/dims=2L	661752 ns	661834 ns	1.00
array/accumulate/Int64/1d	118225 ns	117985 ns	1.00
array/accumulate/Int64/dims=1	75602 ns	75895 ns	1.00
array/accumulate/Int64/dims=1L	1700833 ns	1701305 ns	1.00
array/accumulate/Int64/dims=2	149443 ns	149248 ns	1.00
array/accumulate/Int64/dims=2L	987132 ns	987601 ns	1.00
array/broadcast	18669 ns	18322 ns	1.02
array/broadcast launch	9802 ns	9277 ns	1.06
array/construct	921.8928571428571 ns	894.0714285714286 ns	1.03
array/copy	16422 ns	16295 ns	1.01
array/copyto!/cpu_to_gpu	206280 ns	207658 ns	0.99
array/copyto!/gpu_to_cpu	240748 ns	238782 ns	1.01
array/copyto!/gpu_to_gpu	8825.333333333334 ns	8751.333333333334 ns	1.01
array/iteration/findall/bool	132421 ns	132630 ns	1.00
array/iteration/findall/int	147958 ns	146831 ns	1.01
array/iteration/findfirst/bool	70047 ns	69616 ns	1.01
array/iteration/findfirst/int	71388 ns	71251 ns	1.00
array/iteration/findmin/1d	67493 ns	65622 ns	1.03
array/iteration/findmin/2d	101138 ns	100800 ns	1.00
array/iteration/logical	192381 ns	191177 ns	1.01
array/iteration/scalar	63636 ns	61469 ns	1.04
array/permutedims/2d	49397 ns	48796 ns	1.01
array/permutedims/3d	50739 ns	49965 ns	1.02
array/permutedims/4d	50550 ns	50168 ns	1.01
array/random/rand/Float32	11373 ns	11065 ns	1.03
array/random/rand/Int64	21898 ns	20906 ns	1.05
array/random/rand!/Float32	7540.25 ns	7622.5 ns	0.99
array/random/rand!/Int64	19248 ns	18758 ns	1.03
array/random/randn/Float32	33812 ns	33025 ns	1.02
array/random/randn!/Float32	28396 ns	22888 ns	1.24
array/reductions/mapreduce/Float32/1d	33160 ns	32199 ns	1.03
array/reductions/mapreduce/Float32/dims=1	38014 ns	37930 ns	1.00
array/reductions/mapreduce/Float32/dims=1L	51606 ns	51206 ns	1.01
array/reductions/mapreduce/Float32/dims=2	56431 ns	56676 ns	1.00
array/reductions/mapreduce/Float32/dims=2L	67580 ns	67441 ns	1.00
array/reductions/mapreduce/Int64/1d	41117 ns	40162 ns	1.02
array/reductions/mapreduce/Int64/dims=1	40511 ns	40304 ns	1.01
array/reductions/mapreduce/Int64/dims=1L	88851 ns	88769 ns	1.00
array/reductions/mapreduce/Int64/dims=2	58102 ns	58355 ns	1.00
array/reductions/mapreduce/Int64/dims=2L	83928 ns	83827 ns	1.00
array/reductions/reduce/Float32/1d	33446 ns	32988 ns	1.01
array/reductions/reduce/Float32/dims=1	38019 ns	37878 ns	1.00
array/reductions/reduce/Float32/dims=1L	50983 ns	50721 ns	1.01
array/reductions/reduce/Float32/dims=2	55747 ns	55660 ns	1.00
array/reductions/reduce/Float32/dims=2L	68172 ns	67890 ns	1.00
array/reductions/reduce/Int64/1d	41340 ns	40427 ns	1.02
array/reductions/reduce/Int64/dims=1	40594 ns	40828 ns	0.99
array/reductions/reduce/Int64/dims=1L	88689 ns	88921 ns	1.00
array/reductions/reduce/Int64/dims=2	58258 ns	58360 ns	1.00
array/reductions/reduce/Int64/dims=2L	83867 ns	83383 ns	1.01
array/reverse/1d	17093 ns	17136 ns	1.00
array/reverse/1dL	69886 ns	69806 ns	1.00
array/reverse/1dL_inplace	67702 ns	67603 ns	1.00
array/reverse/1d_inplace	8762.666666666666 ns	8717.666666666666 ns	1.01
array/reverse/2d	20097 ns	20168 ns	1.00
array/reverse/2dL	73682 ns	73471 ns	1.00
array/reverse/2dL_inplace	67534 ns	67397 ns	1.00
array/reverse/2d_inplace	10196 ns	9986 ns	1.02
array/sorting/1d	2643945 ns	2637285 ns	1.00
array/sorting/2d	1018094 ns	1017842 ns	1.00
array/sorting/by	3172488 ns	3173251 ns	1.00
cuda/synchronization/context/auto	1024.6 ns	1004.7 ns	1.02
cuda/synchronization/context/blocking	785.0588235294117 ns	780.5188679245283 ns	1.01
cuda/synchronization/context/nonblocking	5808.6 ns	5923.166666666667 ns	0.98
cuda/synchronization/stream/auto	884.94 ns	847.65 ns	1.04
cuda/synchronization/stream/blocking	666.0445859872611 ns	662.7032258064517 ns	1.01
cuda/synchronization/stream/nonblocking	5739.5 ns	5626 ns	1.02
integration/byval/reference	148183 ns	148148 ns	1.00
integration/byval/slices=1	149278 ns	148882 ns	1.00
integration/byval/slices=2	292264 ns	292021 ns	1.00
integration/byval/slices=3	435373 ns	434930 ns	1.00
integration/cudadevrt	105298 ns	105350 ns	1.00
integration/volumerhs	9136863 ns	9150991 ns	1.00
kernel/indexing	13245 ns	13331 ns	0.99
kernel/indexing_checked	13787 ns	14055 ns	0.98
kernel/launch	2380.3333333333335 ns	2299.8888888888887 ns	1.03
kernel/occupancy	913.0555555555555 ns	883.015873015873 ns	1.03
kernel/rand	17760 ns	17671 ns	1.01
latency/import	4216336889 ns	4202904763 ns	1.00
latency/precompile	4976354862 ns	4971747850 ns	1.00
latency/ttfp	4672194703 ns	4677246466 ns	1.00

This comment was automatically generated by workflow using github-action-benchmark.