Skip to content

compiler: Fix performance profiling of busy-waits - #3055

Open
FabioLuporini wants to merge 6 commits into
mainfrom
fix-wait-timing
Open

FabioLuporini wants to merge 6 commits into
mainfrom
fix-wait-timing

Conversation

@FabioLuporini

Copy link
Copy Markdown
Contributor

No description provided.

@FabioLuporini FabioLuporini self-assigned this Oct 14, 2026
@FabioLuporini FabioLuporini added the bug-C-minor bug in the generated code not affecting correctness label Oct 14, 2026
if size == 1:
continue
if v.is_double_buffering:
# E.g., `f[time] = b[t0]` must become `b1[t0] = b[t0]`

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

actually this is a bug fix (tested both here and in PRO)

@codecov

codecov Bot commented Oct 14, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 98.34711% with 4 lines in your changes missing coverage. Please review.
✅ Project coverage is 84.33%. Comparing base (d7f53a3) to head (8d1a1fa).
⚠️ Report is 2 commits behind head on main.

Files with missing lines Patch % Lines
devito/passes/clusters/buffering.py 55.55% 3 Missing and 1 partial ⚠️
Additional details and impacted files
@@            Coverage Diff             @@
##             main    #3055      +/-   ##
==========================================
+ Coverage   84.23%   84.33%   +0.09%     
==========================================
  Files         258      258              
  Lines       57070    57205     +135     
  Branches     4874     4893      +19     
==========================================
+ Hits        48074    48244     +170     
+ Misses       8169     8134      -35     
  Partials      827      827              
Flag Coverage Δ
pytest-gpu-aomp-amdgpuX 68.60% <70.41%> (+0.10%) ⬆️
pytest-gpu-gcc- 78.93% <22.72%> (-0.08%) ⬇️
pytest-gpu-icx- 78.87% <22.72%> (-0.06%) ⬇️
pytest-gpu-nvc-nvidiaX 69.28% <88.31%> (+0.19%) ⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ 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.

FabioLuporini and others added 5 commits October 19, 2026 13:17
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.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bug-C-minor bug in the generated code not affecting correctness

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants