portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2026-08-14 06:29:36 +01:00
Code Releases Activity
8,304 Commits 95 Branches 14 Tags
6b2ad3db802ea0a2faa755bd919d282ec48e469e
Commit Graph
8304 Commits
This Branch
This Branch
All Branches
Author SHA1 Message Date
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 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 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 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 BoyleandClaude 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 d6b1388741 Modified repack 2026-05-21 12:34:30 -04:00