Barrel shift -> all to all (x2) and distributed FFT work fully load balanced without redundant work.
There is little more I can do now on FFT. Comms dominated and running distributed work dividing bandwidth optimal RingAllToAll
/ccs/home/paboyle/ParallelIO/systems/Frontier/tests/core/Test_fft_prop --mpi 3.6.4.4 --grid 48.48.48.96 --accelerator-threads 8 --shm 4096 --shm-mpi 1 --device-mem
32000 --log Error,Warning,Message,Performance
*************************************************
Benchmarking FFT of LatticeFermionD on plane wave
*************************************************
Grid : Performance : 0.501524 s : FFT took 0.001311 s (transpose P=3)
Grid : Performance : 0.501531 s : FFT pack 5.9e-05 s
Grid : Performance : 0.501533 s : FFT alltoall 0.000828 s
Grid : Performance : 0.501534 s : FFT reorder 0.000204 s
Grid : Performance : 0.501535 s : FFT kernels 1e-05 s
Grid : Performance : 0.501536 s : FFT unpack 5e-05 s
Grid : Performance : 0.509992 s : FFT took 0.001829 s (transpose P=6)
Grid : Performance : 0.510000 s : FFT pack 6e-05 s
Grid : Performance : 0.510002 s : FFT alltoall 0.001436 s
Grid : Performance : 0.510003 s : FFT reorder 0.000202 s
Grid : Performance : 0.510005 s : FFT kernels 9e-06 s
Grid : Performance : 0.510006 s : FFT unpack 5.3e-05 s
Grid : Performance : 0.517690 s : FFT took 0.001599 s (transpose P=4)
Grid : Performance : 0.517698 s : FFT pack 6e-05 s
Grid : Performance : 0.517700 s : FFT alltoall 0.001258 s
Grid : Performance : 0.517701 s : FFT reorder 0.0002 s
Grid : Performance : 0.517702 s : FFT kernels 9e-06 s
Grid : Performance : 0.517703 s : FFT unpack 4.9e-05 s
Grid : Performance : 0.524858 s : FFT took 0.001561 s (transpose P=4)
Grid : Performance : 0.524865 s : FFT pack 5.8e-05 s
Grid : Performance : 0.524867 s : FFT alltoall 0.001213 s
Grid : Performance : 0.524868 s : FFT reorder 0.000209 s
Grid : Performance : 0.524869 s : FFT kernels 8e-06 s
Grid : Performance : 0.524870 s : FFT unpack 4.9e-05 s
*************************************************
FFT of [48 48 48 96] LatticeFermionD took 0.030916 s
*************************************************
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>
Remove the CUB/hipCUB direction entirely. Restore Lattice_reduction_gpu.h,
Lattice_reduction_sycl.h, and Lattice_reduction.h to the state before the
CUB rewrite (commit 969b0a39), recovering the original primary function names
(sumD_gpu_small, sumD_gpu_large, sumD_gpu, sum_gpu, sum_gpu_large) and the
hand-rolled shared-memory reduction kernel.
Delete Lattice_reduction_gpu_cub.h. Update Test_reduction to remove the
old/new comparison sections that depended on sum_gpu_old.
The lesson: CUB DeviceReduce is slower than the hand-rolled kernel for small
types, and the smem sizing problem for the extraction pass has no clean
solution within the accelerator_for abstraction. The right improvement is
a higher radix (12 then 4) in sumD_gpu_large, applied directly to the
existing hand-rolled kernel.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
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>
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>
Peter Boyle
