Peter Boyle
fb5662a449
Setup on Frontier
2026-08-11 16:35:42 -04:00
Peter Boyle
6b2ad3db80
Inverse improvements. Need to make this properly recursive
2026-08-11 16:35:42 -04:00
Peter Boyle
f1a969f0c3
Evict all optoin
2026-08-11 16:35:42 -04:00
Peter Boyle
f18320a152
GeneralCoarsenedMatrix speed up -- same optimisations for fast
...
calculation used by GeneralCoarsenedMatrixMultiRHS
2026-08-11 16:35:42 -04:00
Peter Boyle
bfcde59199
Memory manager improvements
2026-08-11 16:35:42 -04:00
Peter Boyle
f23e9e70cd
Global orthogonalise routine
2026-08-11 16:35:42 -04:00
Peter Boyle
06fcd31da0
Verbose options
2026-08-11 16:35:42 -04:00
Peter Boyle
7e4fe99b1e
Logging update
2026-08-11 16:35:42 -04:00
Peter Boyle
159ea3d64d
Better logging performance
2026-08-11 16:35:42 -04:00
Peter Boyle
1f55c13bc4
Test script for Frontier on I/O
2026-08-11 13:37:50 -04:00
Peter Boyle
d68b111d06
Switch on the faster IO be default
2026-08-11 13:36:50 -04:00
Peter Boyle
41f5a02204
IO updates for AllToAllV aggregation of contiguous chunks.
...
May finally clean up the poor MPI2 IO performance issue that has been persistent.
2026-08-11 13:33:04 -04:00
Peter Boyle
7647576863
Sourceme's for frontier
2026-06-26 11:50:44 -04:00
Peter Boyle
f11ba18df2
Remove accelerator_inline on CPU only code
2026-06-17 20:47:15 +02:00
Peter Boyle
cf8587e401
Update booster compiule
2026-06-17 20:46:53 +02:00
Peter Boyle
7dd35ef749
Make disable accelerator aware mpi compile for CUDA
2026-06-17 20:46:14 +02:00
Peter Boyle
41e570ddce
Annoying old CPU perfmon code should be removed or deprecated as not
...
worth maintaining
2026-06-17 20:45:32 +02:00
Peter Boyle
a452131b50
Print improvement
2026-06-17 20:45:04 +02:00
Peter Boyle
4e49ca55ab
Merge branch 'develop' of https://github.com/paboyle/Grid into develop
2026-06-16 11:20:12 -07:00
Peter Boyle
c3f4474401
Adding Mattia's memory leak test
2026-06-16 11:19:36 -07:00
Peter Boyle and GitHub
3d3eff86f3
Modify move assignment operator to be noexcept
...
Add noexcept specifier to move assignment operator.
2026-06-11 09:44:24 -04:00
Peter Boyle and GitHub
fc9f154ac1
Modify Lattice move constructor to be noexcept
...
Add noexcept specifier to move constructor for Lattice class.
2026-06-11 09:40:06 -04:00
Peter Boyle and GitHub
4aa0bca4dc
Change sum operation to use gpucub mistake in PR from Chris
...
Updated the sum operation definition for GPU reduction to use gpucub instead of cub.
2026-06-01 14:12:25 -04:00
Peter Boyle
905da6f083
Merge branch 'feature/reduction-reorganisation' into develop
2026-05-27 21:01:30 -04:00
Peter Boyle
86c7f29183
Config command update
2026-05-27 16:19:33 -04:00
Peter Boyle
b0c99f876e
Configure on mac update
2026-05-27 16:16:55 -04:00
Peter Boyle
bf5fcdc860
Ease of use for std::complex interchangable with thrust
2026-05-27 16:05:37 -04:00
Peter Boyle
b58a1508fa
Perlmutter cuda version update
2026-05-21 13:25:13 -07:00
Peter Boyle
4d527e81fa
Remove hip specific files
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
7803580aa6
Lattice_reduction_gpu: demote timing logs to Debug, disable by default
...
skills/mpi-heterogeneous: add Bug Class 4 for Frontier GTL/libamdhip64 ABI mismatch
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
32654db366
Test_planned_fft: fix PlannedFFT template parameter to use ::vector_object
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
cd340cfab3
tests: add Test_planned_fft exercising PlannedFFT<vobj>
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
f32866b2ff
tests/fft: remove PlanDestroy calls (FFT handles plans per-call)
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
1cd1dc091e
FFT: add FFTbase, PlannedFFT; factor FFT_dim_execute free function
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
0493656e86
debug: add Test_hipfft_repro — reproducer for hipFFT PARSE_ERROR on ROCm 7
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
66fd504c4d
tests/debug: add G=4 to hipfft fail reproducer
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
be4dd2b52f
tests/debug: test hipMemset variant before cache is populated
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
707d059766
tests/debug: extend hipfft fail reproducer with hipMemset and sync variants
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
f08c755ae6
FFT: use host stack buffer in PlanCreate, not deviceVector
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
dbbfdd4e4b
tests/debug: add minimal hipfft ordering bug fail/pass pair
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
f967fb40bf
tests/debug: test plan-before-malloc vs malloc-before-plan ordering
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
74e0f846cb
tests/debug: extend hipfft reproducer with Grid-realistic howmany and exec tests
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
303a4d26e5
tests/debug: add minimal hipfft plan-creation reproducer
...
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle
119888653c
FFT HIP: use hipfftCreate+hipfftMakePlanMany instead of hipfftPlanMany
2026-05-21 12:34:30 -04:00
Peter Boyle
a9f42c08f9
FFT: pass nullptr for inembed/onembed in hipfftPlanMany to avoid HIPFFT_PARSE_ERROR
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
e79adc9d31
FFT: cache plans per vobj type across calls
...
Plans are created lazily on the first FFT_dim call and reused for all
subsequent calls on the same FFT object. PlanCreate<vobj>() can be
called explicitly to pre-warm the cache. PlanDestroy() must be called
before switching to a different vobj type; the destructor cleans up any
live plans automatically.
Update Test_fft.cc and Test_fftf.cc to call PlanDestroy() between the
LatticeComplex and LatticeSpinMatrix sections that reuse the same FFT object.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
5a9056cd93
Accelerator: lower default accelerator_threads from 16 to 8
...
Benchmark_dwf_fp32 on MI250X GCD: 1.7 TF/s at nt=8, ~300 GF/s at nt=16.
With Nsimd=8 (fp32, GEN_SIMD_WIDTH=64B), nt=8 gives exactly 64 threads =
one full AMD wavefront. Higher values double register demand per block and
hit a register-pressure cliff for stencil kernels.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle
012c36ab5a
Accelerator: raise default accelerator_threads from 2 to 16
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
5c4574f9aa
skills: add gpu-memory-performance.md
...
Documents the acceleratorThreads() default=2 trap, LambdaApply thread
mapping, coalescedRead/Write idiom, when to use __global__ vs
accelerator_for, and fused vs staged HBM access patterns.
Includes observed MI250X numbers from LatticePropagatorD reduction
(50 → 297 → 546 GB/s progression).
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle and Claude Sonnet 4.6
a424775884
sumD_gpu_reduce_words: fuse pack+reduce into single packReduceKernel
...
Replace the two-kernel pack+reduce sequence with a single fused kernel
packReduceKernel<R> that reads R words of each vobj at offset 'base'
and accumulates directly into iVector<iScalar<scalarD>,R>, eliminating
the intermediate bundle buffer entirely.
HBM access per word-group drops from 3x (pack-read + pack-write +
reduce-read) to 1x. Thread count comes from getNumBlocksAndThreads
(warpSize..256) rather than acceleratorThreads(), so occupancy is
correct regardless of the --accelerator-threads setting.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com >
2026-05-21 12:34:30 -04:00
Peter Boyle
Peter BoyleandGitHub
