Repository navigation
Generate !invarant loads instead of ldg calls - #3094
Merged
Merged
Conversation
LLVM 20 removed the llvm.nvvm.ldg.global.{i,f,p} intrinsics in favor
of a regular load from addrspace(1) carrying !invariant.load metadata,
which the NVPTX backend lowers to ld.global.nc (llvm/llvm-project#112834).
Emit that pattern directly via Base.llvmcall on LLVM 20+, while keeping
the old @typed_ccall path for LLVM <20 to preserve Julia 1.10 support.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## master #3094 +/- ##
==========================================
- Coverage 90.43% 90.42% -0.01%
==========================================
Files 141 141
Lines 12025 12025
==========================================
- Hits 10875 10874 -1
- Misses 1150 1151 +1 ☔ View full report in Codecov by Sentry. 🚀 New features to boost your workflow:
|
Contributor
There was a problem hiding this comment.
CUDA.jl Benchmarks
Details
| Benchmark suite | Current: aa2b581 | Previous: 33ffdef | Ratio |
|---|---|---|---|
array/accumulate/Float32/1d |
100543.5 ns |
101156 ns |
0.99 |
array/accumulate/Float32/dims=1 |
75913 ns |
76494 ns |
0.99 |
array/accumulate/Float32/dims=1L |
1584818 ns |
1583795 ns |
1.00 |
array/accumulate/Float32/dims=2 |
142546 ns |
143348 ns |
0.99 |
array/accumulate/Float32/dims=2L |
657124.5 ns |
656949 ns |
1.00 |
array/accumulate/Int64/1d |
118044 ns |
118389 ns |
1.00 |
array/accumulate/Int64/dims=1 |
79608 ns |
79916 ns |
1.00 |
array/accumulate/Int64/dims=1L |
1694179 ns |
1694529 ns |
1.00 |
array/accumulate/Int64/dims=2 |
154776 ns |
155647 ns |
0.99 |
array/accumulate/Int64/dims=2L |
961049 ns |
961265 ns |
1.00 |
array/broadcast |
20728.5 ns |
20513 ns |
1.01 |
array/construct |
1365.5 ns |
1333.8 ns |
1.02 |
array/copy |
18713 ns |
18686 ns |
1.00 |
array/copyto!/cpu_to_gpu |
212771 ns |
212224.5 ns |
1.00 |
array/copyto!/gpu_to_cpu |
281304 ns |
282985 ns |
0.99 |
array/copyto!/gpu_to_gpu |
11352 ns |
11526 ns |
0.98 |
array/iteration/findall/bool |
131491 ns |
131867 ns |
1.00 |
array/iteration/findall/int |
148274 ns |
148383 ns |
1.00 |
array/iteration/findfirst/bool |
81473 ns |
81401 ns |
1.00 |
array/iteration/findfirst/int |
83566 ns |
83519 ns |
1.00 |
array/iteration/findmin/1d |
86140.5 ns |
83773.5 ns |
1.03 |
array/iteration/findmin/2d |
116882 ns |
117018 ns |
1.00 |
array/iteration/logical |
198710.5 ns |
197560.5 ns |
1.01 |
array/iteration/scalar |
66988 ns |
68375 ns |
0.98 |
array/permutedims/2d |
52061 ns |
52136 ns |
1.00 |
array/permutedims/3d |
52846 ns |
52817.5 ns |
1.00 |
array/permutedims/4d |
51400 ns |
52006 ns |
0.99 |
array/random/rand/Float32 |
12902 ns |
13083 ns |
0.99 |
array/random/rand/Int64 |
29429 ns |
37254 ns |
0.79 |
array/random/rand!/Float32 |
8553 ns |
8673.666666666666 ns |
0.99 |
array/random/rand!/Int64 |
34011 ns |
33174 ns |
1.03 |
array/random/randn/Float32 |
37353 ns |
43918 ns |
0.85 |
array/random/randn!/Float32 |
31240 ns |
31437 ns |
0.99 |
array/reductions/mapreduce/Float32/1d |
34690 ns |
35490 ns |
0.98 |
array/reductions/mapreduce/Float32/dims=1 |
49221 ns |
45840 ns |
1.07 |
array/reductions/mapreduce/Float32/dims=1L |
51519 ns |
51786 ns |
0.99 |
array/reductions/mapreduce/Float32/dims=2 |
56514 ns |
56654 ns |
1.00 |
array/reductions/mapreduce/Float32/dims=2L |
69396 ns |
69672 ns |
1.00 |
array/reductions/mapreduce/Int64/1d |
42470 ns |
43204 ns |
0.98 |
array/reductions/mapreduce/Int64/dims=1 |
42017 ns |
42387 ns |
0.99 |
array/reductions/mapreduce/Int64/dims=1L |
87541 ns |
87636 ns |
1.00 |
array/reductions/mapreduce/Int64/dims=2 |
59721 ns |
59567 ns |
1.00 |
array/reductions/mapreduce/Int64/dims=2L |
84874 ns |
85154 ns |
1.00 |
array/reductions/reduce/Float32/1d |
34722 ns |
35536 ns |
0.98 |
array/reductions/reduce/Float32/dims=1 |
49278 ns |
49541 ns |
0.99 |
array/reductions/reduce/Float32/dims=1L |
51918 ns |
51867 ns |
1.00 |
array/reductions/reduce/Float32/dims=2 |
56546 ns |
56724 ns |
1.00 |
array/reductions/reduce/Float32/dims=2L |
69660 ns |
70152 ns |
0.99 |
array/reductions/reduce/Int64/1d |
42651 ns |
43180 ns |
0.99 |
array/reductions/reduce/Int64/dims=1 |
43921.5 ns |
46092.5 ns |
0.95 |
array/reductions/reduce/Int64/dims=1L |
87453 ns |
87572 ns |
1.00 |
array/reductions/reduce/Int64/dims=2 |
59400 ns |
59560 ns |
1.00 |
array/reductions/reduce/Int64/dims=2L |
84641.5 ns |
84946 ns |
1.00 |
array/reverse/1d |
18234 ns |
18382 ns |
0.99 |
array/reverse/1dL |
68839 ns |
68959.5 ns |
1.00 |
array/reverse/1dL_inplace |
65874 ns |
65940 ns |
1.00 |
array/reverse/1d_inplace |
8550 ns |
8534 ns |
1.00 |
array/reverse/2d |
20657.5 ns |
20498 ns |
1.01 |
array/reverse/2dL |
73218.5 ns |
72575.5 ns |
1.01 |
array/reverse/2dL_inplace |
65983 ns |
66219 ns |
1.00 |
array/reverse/2d_inplace |
10137 ns |
10379 ns |
0.98 |
array/sorting/1d |
2735364 ns |
2734628.5 ns |
1.00 |
array/sorting/2d |
1068300 ns |
1068430 ns |
1.00 |
array/sorting/by |
3303721 ns |
3303603 ns |
1.00 |
cuda/synchronization/context/auto |
1168 ns |
1179.7 ns |
0.99 |
cuda/synchronization/context/blocking |
936.1607142857143 ns |
917.5135135135135 ns |
1.02 |
cuda/synchronization/context/nonblocking |
7503.4 ns |
7659 ns |
0.98 |
cuda/synchronization/stream/auto |
988.8823529411765 ns |
1056.923076923077 ns |
0.94 |
cuda/synchronization/stream/blocking |
824.1047619047619 ns |
797.961038961039 ns |
1.03 |
cuda/synchronization/stream/nonblocking |
7081.299999999999 ns |
8313.1 ns |
0.85 |
integration/byval/reference |
144087 ns |
144149 ns |
1.00 |
integration/byval/slices=1 |
145751 ns |
146090 ns |
1.00 |
integration/byval/slices=2 |
284544 ns |
284910 ns |
1.00 |
integration/byval/slices=3 |
423048 ns |
423319 ns |
1.00 |
integration/cudadevrt |
102570 ns |
102707 ns |
1.00 |
integration/volumerhs |
9435863 ns |
9433360.5 ns |
1.00 |
kernel/indexing |
13307 ns |
13349 ns |
1.00 |
kernel/indexing_checked |
14030 ns |
14195 ns |
0.99 |
kernel/launch |
2128.3333333333335 ns |
2182.222222222222 ns |
0.98 |
kernel/occupancy |
683.0331125827814 ns |
661.179012345679 ns |
1.03 |
kernel/rand |
17768 ns |
14520 ns |
1.22 |
latency/import |
3808580304.5 ns |
3811597161 ns |
1.00 |
latency/precompile |
4588586477 ns |
4611897362.5 ns |
0.99 |
latency/ttfp |
4426087177 ns |
4392419951 ns |
1.01 |
This comment was automatically generated by workflow using github-action-benchmark.
kshyatt
pushed a commit
that referenced
this pull request
Aug 5, 2026
LLVM 20 removed the llvm.nvvm.ldg.global.{i,f,p} intrinsics in favor
of a regular load from addrspace(1) carrying !invariant.load metadata,
which the NVPTX backend lowers to ld.global.nc (llvm/llvm-project#112834).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Closes #2531
Review with whitespace hidden