portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2026-08-28 13:29:35 +01:00
Code Releases Activity
8,354 Commits 95 Branches 14 Tags
e9be1a91b35a27aed1e91058d16932dbf788fe57
Commit Graph
100 Commits
This Branch
This Branch
All Branches
Author SHA1 Message Date
Peter Boyle e9be1a91b3 Speed up on mac arm with --enable-simd=NEONv8.3 2026-08-19 19:49:59 -04:00
Peter Boyle 7db3b96d2d Deprecate D2 fields 2026-08-19 19:49:26 -04:00
Peter Boyle 429ad2a71a sComplex and lexLattice prep work 2026-08-19 19:42:56 -04:00
Peter Boyle 0b1598e9ef Annoying warn on printf fix 2026-08-19 19:41:43 -04:00
Peter Boyle 8d04b91932 Clean up whitespace. Gauge group needs a Gimpl or some other policy parameter. 
Perhaps rename iGaugeGroupd and pass EITHER vComplexD, vComplexF OR sCompexD / sComplexF
Use a typedef to restore the original Gauge group
 2026-08-19 19:39:42 -04:00
Peter Boyle 665094d94b Lex gauge actions. 2026-08-19 19:38:48 -04:00
Peter Boyle 919d126af1 Remove D2 lattice fields; remove WilsonCompressor based half precision project; add sComplex lexLattice layouts 
Fix staggered instantiations for new types
 2026-08-19 19:37:29 -04:00
Peter Boyle abd5df1729 Discourage use of D2 lattices -- it was low performance. 
Introduce scalar lexLattice layout
 2026-08-19 19:35:59 -04:00
Peter Boyle 3243ad91e8 Fix race condition during read after write due to missing barrier. 
Enable both lexLattice (scalar order) and Lattice (vector or partial SOA order)
 2026-08-19 19:34:54 -04:00
Peter Boyle c10d451f4d Enable interafaces with both coarse and fine being either single higher dim multiRHS lattice OR vector of lower dim fields 2026-08-19 19:33:55 -04:00
Peter Boyle 9c4a61850e Features for better BLAS based coarse grids, for both single and multi RHS. 
Use sComplex and lexLattice in future to avoid layout changes and copies
 2026-08-19 19:32:36 -04:00
Peter Boyle 5e3fe24412 Recursive Schur complement based distributed matrix inverse by Claude after discussing plan and instructions from me 2026-08-19 19:31:46 -04:00
Peter Boyle 7b81ca6b52 Changes to make Nsimd=1 sComplex lexLattices functional for "where" statements 2026-08-19 19:30:34 -04:00
Peter Boyle 8c9ea4cde2 Json "" warning remove 2026-08-19 19:29:27 -04:00
Peter Boyle 0ac72cb6a3 Don't use host as it sucks. 2026-08-15 12:24:49 -04:00
Peter Boyle c4b089cf15 More timing updates 2026-08-15 09:51:56 -04:00
Peter Boyle 85e1bbf4bc Added gather benchmark to cross refer to the 'lazy' analogue of summing a vector with many zeroes in regions not owned by THIS rank 2026-08-15 09:51:12 -04:00
Peter Boyle 032f111c8b Instrument as not faster 2026-08-15 07:42:18 -04:00
Peter Boyle e38528888a Submission scripts for distributed schur coarse grid inverse 2026-08-14 18:28:46 -04:00
Peter Boyle 7c9a6d354e Updates for thrust 2026-08-14 18:21:45 -04:00
Peter Boyle 3ec90803ac no compil patches 2026-08-14 18:21:13 -04:00
Peter Boyle c22473f15d First test of distributed schur recursive inverse 2026-08-14 17:53:35 -04:00
Peter Boyle 55c064de28 Batch blas test extension 2026-08-14 17:53:35 -04:00
Peter Boyle 06ce057920 Bug fix for CPU build 2026-08-14 17:53:35 -04:00
Peter Boyle 76b4bd6d12 Distributed inverse support 2026-08-14 17:53:35 -04:00
Peter Boyle b5541aab55 Schur Inverse support 2026-08-14 17:53:35 -04:00
Peter Boyle a7160ac513 Support for distributed Schur inverse 2026-08-14 17:53:35 -04:00
Peter Boyle 02d0301c9f Scaling perf 2026-08-13 23:46:33 -04:00
Peter Boyle a6cdf20c18 IO tests 2026-08-13 20:08:56 -04:00
Peter Boyle ad9a413892 Update 2026-08-13 20:08:25 -04:00
Peter Boyle 1fddd2c29b Updates to IO regression example for MPICH issue 2026-08-13 18:54:27 -04:00
Peter Boyle 2f75067569 Some I/O simple benchmarks 2026-08-13 18:54:26 -04:00
Peter Boyle 6e8a00f215 Simplest starting point 2026-08-13 16:57:10 -04:00
Peter Boyle 702773e5fb BatchedBlas.h updates. May cause some churn. Trying to avoid scalar pointer copies to device. 2026-08-13 14:55:03 -04:00
Peter Boyle 4dfbd850ff Buf fix on DenseCoarseMatrix.h, coefficient caching optimisation on BatchGEMM 2026-08-13 13:26:42 -04:00
Peter Boyle b039e659af 4D pseudo fermion 2026-08-12 12:51:55 -04:00
Peter Boyle d16d44dda0 Prepare for reduced Ls during evolution forces using 4D pseudo fermion. 
Worth an effort to understand if the different stochastic estimator is cleaner.
 2026-08-12 12:51:01 -04:00
Peter Boyle 1c19389ba6 Two flavour boson term 2026-08-12 12:50:32 -04:00
Peter Boyle 02fdff674c Test the PVdagM left prec pseudofermion formulation 2026-08-12 12:49:59 -04:00
Peter Boyle fd8b6a23a6 PvDagM or other left prec precon two flavour ratio 2026-08-12 12:49:21 -04:00
Peter Boyle 9e3a51d078 Use the new DenseCoarseInverse class 2026-08-12 12:48:19 -04:00
Peter Boyle 6f7a2ad7c7 Dense coarse inverse wrapper class, bundling up a bunch of useful work 2026-08-12 12:47:40 -04:00
Peter Boyle 499d656949 Batch blas test 2026-08-12 12:47:16 -04:00
Peter Boyle ba68f09026 Batch inverse 2026-08-12 12:46:56 -04:00
Peter Boyle 3bdeeb73ef Update 2026-08-12 12:44:56 -04:00
Peter Boyle 19868a800f Blas optimise change 2026-08-12 12:43:40 -04:00
Peter Boyle df908ee872 3 level dense 2026-08-11 17:50:12 -04:00
Peter Boyle 84715ff4b9 Default to fast path 2026-08-11 17:39:12 -04:00
Peter Boyle 5792195073 Check in codes used at time of Lattice conference for PVdagM multigrid 2026-08-11 16:35:42 -04:00
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 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
Peter BoyleandClaude Sonnet 4.6 796c6cae4e Enable GRID_REDUCTION_TIMING unconditionally 
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-21 12:34:30 -04:00
Peter BoyleandClaude Sonnet 4.6 1a8064d6d9 Lattice_reduction_gpu: add GRID_REDUCTION_TIMING instrumentation 
Uncomment #define GRID_REDUCTION_TIMING to enable per-phase timing output:

  sumD_gpu_reduce_words: pack time (accelerator_for) per R and base
  sumD_gpu_small:        reduceKernel+barrier time and D2H time separately
  sumD_gpu_large:        total wall time across all word groups

This lets us identify whether the large-type bottleneck is in the pack
kernel, the shared-memory reduction kernel, the barrier, or the D2H.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-21 12:34:30 -04:00
Peter BoyleandClaude Sonnet 4.6 43648924c3 sumD_gpu_large: radix-12 word-bundle reduction replacing radix-1 
Replace the word-by-word loop (one kernel launch per scalar word) with
sumD_gpu_reduce_words<R> which packs R consecutive vector_type words per
site into iVector<iScalar<vector>,R>, then calls the existing sumD_gpu_small
shared-memory kernel once for the whole bundle.

Dispatch: radix-12 first, radix-4 for the remainder < 12, radix-1 for
any final < 4 words.  For LatticePropagator (144 words = 12x12), this
reduces the kernel-launch count from 144 to 12 -- a 12x reduction.

Bundle::Nsimd() inherits from vector_type so sumD_gpu_small handles SIMD
lane extraction and double-precision promotion identically to the scalar
word case.  sizeof(Bundle::scalar_objectD) = R*16 <= 192 B; well within
sharedMemPerBlock on all supported devices.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-21 12:34:30 -04:00