Repository navigation
compiler: Fix performance profiling of busy-waits - #3055
Merged
Merged
Conversation
FabioLuporini
added
the
bug-C-minor
FabioLuporini
commented
Oct 14, 2026
| if size == 1: | ||
| continue | ||
| if v.is_double_buffering: | ||
| # E.g., `f[time] = b[t0]` must become `b1[t0] = b[t0]` |
Contributor
Author
There was a problem hiding this comment.
actually this is a bug fix (tested both here and in PRO)
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #3055 +/- ##
==========================================
+ Coverage 84.15% 84.25% +0.10%
==========================================
Files 258 258
Lines 57267 57398 +131
Branches 4893 4911 +18
==========================================
+ Hits 48191 48361 +170
+ Misses 8249 8211 -38
+ Partials 827 826 -1
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
mloubout
approved these changes
Oct 14, 2026
FabioLuporini
force-pushed
the
fix-wait-timing
branch
2 times, most recently
from
October 15, 2026 08:39
bd1fd36 to
ea8fc74
Compare
FabioLuporini
force-pushed
the
fix-wait-timing
branch
2 times, most recently
from
October 21, 2026 08:59
8d1a1fa to
d07489c
Compare
Select compute_100a/sm_100a in CudaCompiler for detected SM100 devices. Use explicit Blackwell as the SM100 fallback when device detection is unavailable, and preserve other detected targets. Explicit gencode avoids the generic compute_100 PTX fallback generated by -arch=sm_100a, which rejects architecture-specific instructions such as tcgen05. Target selection no longer needs a downstream pass override. Add hardware-independent coverage of native fallback, older GPUs, SM100, and other Blackwell targets.
Reuse the first buffer's slot indices when inserting further buffering levels, including initialization outside the time loop. Preserve physical irregular-buffer slots in copy-in masks and flatten transfers using their allocated stride independently of the transferred payload size. Add a subsampled multilevel buffering regression for degrees 1, 2 and 4.
The lock and flag an asynchronous task synchronises on are plain volatile ints. The fences added for them order nothing on their own: a release fence pairs with an acquire fence only through atomic accesses, so a weakly ordered CPU may e.g. see a lock released before the data it guards. Access every lock and flag with __atomic acquire loads and release stores, which are valid in C and C++ and need no declaration change, and drop the now redundant ThreadFence.
FabioLuporini
force-pushed
the
fix-wait-timing
branch
from
October 21, 2026 10:26
d07489c to
b494165
Compare
EdCaunt
reviewed
Oct 21, 2026
| iet, _ = super()._make_waitlock(iet, sync_ops, *args) | ||
| waitloop, *body = iet.body | ||
|
|
||
| # NVHPC 25.7 miscompiles tiled loops following an inline atomic wait |
Contributor
There was a problem hiding this comment.
Is this documented anywhere? Perhaps a link to said documentation wouldn't go amiss?
Contributor
Author
There was a problem hiding this comment.
what documentation? it's an actual NVC compiler bug, like several we had to workaround in the past..
Comment on lines
+75
to
+76
| flag = FieldFromComposite(sdata.symbolic_flag, sdata[d]) | ||
| condition = CondNe(atomic_load(flag), 1) |
Contributor
There was a problem hiding this comment.
Hoist outside the if block?
| (BLACKWELL, 103, '-arch=sm_103'), | ||
| (BLACKWELL, 120, '-arch=sm_120') | ||
| ]) | ||
| def test_cuda_target(monkeypatch, platform, cc, target): |
Contributor
There was a problem hiding this comment.
Is there a better approach than using monkeypatching?
Contributor
Author
There was a problem hiding this comment.
in general probably yes, in this specific case I'm not so sure tbf
mloubout
approved these changes
Oct 21, 2026
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
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.
No description provided.