portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2026-07-17 15:43:27 +01:00
Code Releases Activity
8,195 Commits 95 Branches 14 Tags
baa70d8ec99ade706d73490c05969447569d2a63
Commit Graph
100 Commits
This Branch
This Branch
All Branches
Author SHA1 Message Date
Peter BoyleandClaude Sonnet 4.6 baa70d8ec9 Test_reduction: add timing benchmark for new vs old reduction paths 
Reports us/call and GB/s for sum_gpu (CUB/sycl::reduction) and
sum_gpu_old (hand-rolled shared-memory) for each field type, with
5-call warmup and 100-call timed loop.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-18 12:31:13 -04:00
Peter BoyleandClaude Sonnet 4.6 c93b338bdd skills: HPC battle-hardening skill files for GPU+MPI correctness 
Six skill files encoding expertise for making codebases robust on
problematic HPC systems, covering: correctness verification
(double-run, fingerprinting, flight recorder), hang diagnosis,
GPU runtime correctness (premature barrier, infinite poll),
MPI correctness on heterogeneous systems (device buffer aliasing,
AARCH64 PLT corruption, deterministic reductions),
compiler validation, and communication/computation overlap pipeline
design.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-18 12:10:44 -04:00
Peter BoyleandClaude Sonnet 4.6 c0472aa0ec Test_reduction: use separate float and double grids 
Float fields require a grid constructed with vComplexF::Nsimd(); using
a double grid causes grid->_gsites to undercount the sites in float
vobjF, making the constant-field expected value wrong.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-18 12:09:35 -04:00
Peter BoyleandClaude Sonnet 4.6 09552cfd73 Rename scalarNorm2 to squaredSum in Test_reduction.cc 
The function computes |sum|^2 — the squared magnitude of an aggregate sum —
not a norm. squaredSum makes clear that squaring is applied to the sum, not
to individual site values before summing, distinguishing it from sumOfSquares
(the squared L2 norm).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-15 23:15:11 -04:00
Peter BoyleandClaude Sonnet 4.6 003fec509c Fix Zero() used on thrust::complex in WordBundle4 initialisation 
Grid's Zero() sentinel is not assignable to thrust::complex<double>;
use scalarD(0) instead.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-15 18:10:17 -04:00
Peter BoyleandClaude Sonnet 4.6 773a82d87f Reinstate large/small dispatch in CUB reduction path; radix-4 word-bundle for large types 
rocPRIM's DeviceReduce requires warpSize(64) threads each holding one element in shared
memory, so sizeof(T)*64 must fit in sharedMemPerBlock.  LatticePropagator::scalar_objectD
is 2304 bytes (64*2304 = 147 KB), exceeding the budget and triggering a compile-time
static_assert in limit_block_size.

Introduce sumD_gpu_direct (the original direct-CUB path, safe for small types) and a new
sumD_gpu_large that groups the vobj's vector_type words in bundles of 4, reducing each
bundle as WordBundle4<scalarD> (64 bytes, 64*64 = 4 KB — always within budget).  If
words % 4 != 0, the final partial bundle is zero-padded.  sumD_gpu dispatches at compile
time via if constexpr on sizeof(sobjD) > 512.

For LatticePropagator (144 words) this gives 36 CUB launches instead of 144.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-15 16:55:58 -04:00
Peter BoyleandClaude Sonnet 4.6 286c29d6fb Add Test_reduction to tests/debug 
Tests the new CUB/hipCUB/SYCL lattice reduction (sum_gpu) against the
preserved hand-rolled implementation (sum_gpu_old) for LatticeComplexF/D,
LatticeColourMatrixF/D and LatticePropagatorF/D.

Part a) gaussian random field: checks that old and new agree to within
float/double roundoff tolerance.
Part b) constant field (= 1.0, identity-matrix init): verifies
innerProduct(sum, sum) = Ncomp * V^2 where Ncomp counts the nonzero
diagonal scalar components per site (1 / Nc / Ns*Nc respectively).

Make.inc is auto-generated by scripts/filelist on bootstrap and is not
tracked; the new .cc file is all that is needed.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-15 14:31:33 -04:00
Peter BoyleandClaude Sonnet 4.6 969b0a3922 Rewrite lattice GPU reduction to use CUB, hipCUB, and SYCL reduction 
Replace hand-rolled shared-memory reduction kernels (reduceBlock/reduceBlocks/
reduceKernel) and the global device variable retirementCount with a unified
CUB/hipCUB DeviceReduce::Reduce path for CUDA/HIP and sycl::reduction for SYCL.
No small/large split is needed: both CUB and sycl::reduction handle arbitrary
object sizes internally.

Old implementations preserved as sum_gpu_old / sumD_gpu_old etc. in the
original files for regression testing on GPU hardware.

Also add CLAUDE.md with build, test, and architecture guidance.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-15 13:41:56 -04:00
Peter Boyle c6c2834e03 Hip Happy 2026-05-15 11:30:29 -04:00
Peter Boyle 856545a1db Support ROCM 7.0.2 2026-05-15 11:30:29 -04:00
Peter Boyle b37390bb5a 4 node usqcd run 2026-04-27 14:40:11 -07:00
Peter Boyle 829dc8cceb 32 node 2026-04-27 14:38:02 -07:00
Peter Boyle 13cc2c39f5 FOM run 2026-04-27 14:20:49 -07:00
Peter Boyle 66ea3b271c Merge branch 'develop' of https://github.com/paboyle/Grid into develop 2026-04-27 13:55:52 -07:00
Peter Boyle d293b58a20 384 node baseline run 2026-04-27 13:54:40 -07:00
Peter Boyle ce093b2bf3 rdtsc 2026-04-27 13:54:06 -07:00
Peter Boyle e4404efe5a Perlmutter compile update 2026-04-27 13:53:28 -07:00
Peter Boyle 5ce270f1de Adding Claude related files 2026-04-21 10:41:18 -04:00
Peter Boyle af43b067a0 New CLAUDE controllable visualiser 2026-04-10 11:23:25 -04:00
Peter Boyle 595ceaac37 Include grid header and make the ENABLE correct 2026-03-11 17:24:44 -04:00
Peter Boyle daf5834e8e Fixing incorrect PR about disable fermion instantiations 2026-03-11 17:05:46 -04:00
Peter Boyle 0d8658a039 Optimised 2026-03-05 06:06:32 -05:00
Peter Boyle 095e004d01 Setup change GCR 2026-03-05 06:06:32 -05:00
Peter Boyle 0acabee7f6 Modest change 2026-03-05 06:06:32 -05:00
Peter Boyle 76fbcffb60 Improvement to 16^3 hdcg 2026-03-05 06:06:32 -05:00
Peter Boyle 2a8084d569 Subspace setup 2026-02-13 17:26:11 -05:00
Peter Boyle 6ff29f9d4f Alternate multigrids 2026-02-13 17:25:45 -05:00
Peter Boyle 7cd3f21e6b preserving a bunch of experiments on setup and g5 subspace doubling 2026-01-06 05:57:39 -05:00
Peter Boyle 74e6b19f83 Looks like the reuse of xfers in staggered has bugs or corner cases depending on volume 2025-11-17 22:29:06 -05:00
Peter Boyle 2e684028de Improvements 2025-11-14 18:12:27 -05:00
Peter Boyle fe0db53842 FFT offload to GPU and MUCH faster comms. 
40x speed up on Frontier
 2025-08-21 16:45:38 -04:00
Peter Boyle 76c0ada1e1 Benchmark for En Hung 2025-08-21 16:45:38 -04:00
Peter Boyle 7aa06329d0 Update for new stencil compression options 2025-06-17 18:06:19 +02:00
Peter Boyle 9d6a38c44c Compressed comms options as Sloppy 2025-06-17 16:43:53 +02:00
Peter Boyle 6ec5cee368 Preparing for compressed comms 2025-06-17 16:38:10 +02:00
Peter Boyle f2e9a68825 Simplify 2025-06-13 17:32:05 +02:00
Peter Boyle d88750e6b6 Sloppy + non-sloppy 2025-06-13 16:42:01 +02:00
Peter Boyle 821358eda7 Remove partial dirichlet. Favour intro reduced prec comms options 2025-06-13 05:08:45 +02:00
Peter Boyle fce6e1f135 Kill core files for quota reasons 2025-06-13 05:08:15 +02:00
Peter Boyle 8f0bb3e676 remove partial dirichlet 2025-06-13 05:07:56 +02:00
Peter Boyle 262c70d967 USe sloppy comms options 2025-06-13 05:07:23 +02:00
Peter Boyle da43ef7c2d REmove partial dirichlet option. It's going nowhere 2025-06-13 05:05:15 +02:00
Peter Boyle 7b60ab5df1 Warning suppress 2025-06-13 05:04:55 +02:00
Peter Boyle f6b961a64e Warning suppress 2025-06-13 05:04:47 +02:00
Peter Boyle f1ed988aa3 Interface to reduced precision comms 2025-06-13 05:04:12 +02:00
Peter Boyle eea51bb604 Suppress annoying warns 2025-06-13 05:03:36 +02:00
Peter Boyle 9203126aa5 Scripts 2025-06-11 15:30:16 +02:00
Peter Boyle f90ba4712a Update for Jupiter 2025-06-11 15:24:34 +02:00
Peter Boyle 3737a24096 Updated python output 2025-06-03 14:09:29 -04:00
Peter Boyle dc546aaa4b Updated config options for BNL cluster 2025-05-13 18:44:47 -04:00
Peter Boyle 5364d580c9 Output chirality, eigenvector density files and python source lego plot 2025-05-13 18:44:47 -04:00
Peter Boyle 2a9a6347e3 Do not require Grid format RNGs and also to the 5Li reporting 2025-05-13 18:44:47 -04:00
Peter Boyle cfdb56f314 Run measurements at t=0 too 2025-05-13 18:44:46 -04:00
Peter Boyle d60a80c098 Fixes and visualisation 2025-04-29 18:04:23 -04:00
Peter Boyle bb8b6d9d73 Fix 2025-04-29 18:04:04 -04:00
Peter Boyle 677b4cc5b0 Make all tests compile 2025-04-24 20:33:26 -04:00
Peter Boyle be565ffab6 update mac config command 2025-04-24 14:50:06 -04:00
Peter Boyle df6120e5f6 CPU compile oops fix 2025-04-24 14:50:06 -04:00
Peter Boyle 6d015ae8fc Visualisation tools 2025-04-24 13:47:34 -04:00
Peter Boyle 233150d93f Bug fix for no accelerator aware MPI, thanks Shuhei for finding it. 2025-04-24 11:40:46 -04:00
Peter Boyle 7af8c77a52 Normalise 2025-04-24 11:37:39 -04:00
Peter Boyle e652fc2825 Shared Memory test reenabled on every Grid object creation. 
Const improvements in Accelerator.h
 2025-04-07 11:51:40 -04:00
Peter Boyle a49fa3f8d0 ROCM 6.3.1 appears to work 2025-04-07 11:50:59 -04:00
Peter Boyle cd452a2f91 Slurm update 2025-04-04 18:40:20 -04:00
Peter Boyle 4f89f603ae Changes to add back shared memory test on GPU 2025-04-04 18:40:15 -04:00
Peter Boyle 11dc2c5e1d PVdagM initialise 2025-04-04 18:35:06 -04:00
Peter Boyle 6fec3c15ca Cleaner printing 2025-04-04 18:35:06 -04:00
Peter Boyle 938c47480f Updated compile on frontier. 
Unsatisfactory hacsk
 2025-04-04 18:35:06 -04:00
Peter Boyle 3811d19298 Fence 2025-04-04 18:35:06 -04:00
Peter Boyle 83a3ab6b6f Barrier -- not sure 100% this was needed 2025-04-04 18:35:05 -04:00
Peter Boyle d66a9af6a3 No compile fix 2025-04-04 18:35:05 -04:00
Peter Boyle adc90d3a86 NVLINK GET/PUT on cuda aware mpi 2025-04-04 18:35:05 -04:00
Peter Boyle ebbd015c5c Deprecate shared memory copy as direction matters on nvidia GPU 2025-04-04 18:35:05 -04:00
Peter Boyle 4ab73b36b2 Deprecate shared memory copy as direction matters on GPU 2025-04-04 18:35:05 -04:00
Peter Boyle 130e07a422 Non hermitian support 2025-04-04 18:35:05 -04:00
Peter Boyle 8f47bb367e Shifted non herm 2025-04-04 18:35:05 -04:00
Peter Boyle 0c3cb60135 Script update 2025-04-04 18:35:05 -04:00
Peter Boyle 9eae8fca5d Size outut 2025-04-04 18:35:05 -04:00
Peter Boyle 882a217074 Example of Useful prerequisite installs with spack 2025-03-26 11:28:53 -04:00
Peter Boyle 3d014864e2 Makinig LLVM happy 2025-03-06 14:19:25 -05:00
Peter Boyle c74d11e3d7 PVdagM MG 2025-02-01 11:04:13 -05:00
Peter Boyle 5a4f9bf2e3 Force the ROCM version 2024-10-29 18:12:31 -04:00
Peter Boyle b91fc1b6b4 Merge branch 'feature/boosted' into feature/deprecate-uvm 
Fixed boosted free field test
 2024-10-28 16:53:09 -04:00
Peter Boyle eafc150034 Test fft asserts 2024-10-23 16:46:26 -04:00
Peter Boyle 2877f1a268 Verbose reduce 2024-10-23 15:14:16 -04:00
Peter Boyle 1e893af775 GPU happy 2024-10-23 14:52:15 -04:00
Peter Boyle d9f430a575 Happy GPU 2024-10-23 14:51:16 -04:00
Peter Boyle 63abe87f36 Memory manager verbose improvements that were useful to track an error 2024-10-23 14:49:13 -04:00
Peter Boyle 368d649c8a feature/deprecate-uvm happier -- preallocate device resident neigbour table 2024-10-23 14:47:55 -04:00
Peter Boyle 5603464f39 Fix in partial fraction import/export physical and 
make the GPU happier on the deprecate-uvm -- don't use static vectors, make member of class
 2024-10-23 14:45:58 -04:00
Peter Boyle 655c79f39e Suppress warning on partial override 2024-10-23 14:44:41 -04:00
Peter Boyle 565b231c03 Nvcc happy 2024-10-23 14:44:17 -04:00
Peter Boyle 62a9f180fa NVCC happy 2024-10-23 14:44:04 -04:00
Peter Boyle 5ae77876a8 Meson field and Aslash field on GPU; some compiler warning removed 2024-10-18 19:08:06 -04:00
Peter Boyle 4ed2c2c74f Config command 2024-10-18 13:58:33 -04:00
Peter Boyle 955da582b6 Working on NVCC 2024-10-18 13:58:03 -04:00
Peter Boyle 11b07b950d Vanilla linux compile, assuming spack prerequisites 2024-10-18 13:57:40 -04:00
Peter Boyle 8f70cfeda9 Clean up 2024-10-18 13:56:53 -04:00
Peter Boyle ce64271048 Remove the copying version 2024-10-18 13:56:24 -04:00
Peter Boyle 6815e138b4 Boosted fermion attempt 2024-10-17 18:37:33 +01:00