Skip to content

Commit eef2da0

Browse files
alfCclaude
andcommitted
add FFT size-sweep benchmark vs FFTW with gnuplot output
- benchmark/algorithms_fft.cpp: sweeps 1-D (16..2^21), 2-D (16^2..2048^2), 3-D (8^3..256^3); plan-recycled execution timing, cache flushed per rep, FFTW_ESTIMATE, single thread; writes gnuplot-friendly fft_bench_{1d,2d,3d}.dat (n, N, ms, benchFFT MFLOPS, mine/FFTW ratio) - benchmark/algorithms_fft_plots.gp: renders the three speed-vs-size plots - register algorithms_fft.cpp.x in the standalone benchmark project - fft.NOTES.md: add CUDA adaptation plan (section 8) Measured on this machine: faster than FFTW for 1-D n <= 128, 2-D >= 256^2 (1.4-1.7x), and 3-D 256^3 (2.2x); FFTW leads mid-range 1-D (SIMD codelets). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
1 parent ae1c868 commit eef2da0

4 files changed

Lines changed: 202 additions & 0 deletions

File tree

benchmark/CMakeLists.txt

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,11 @@ target_link_libraries(matrix_multiplication.cpp.x benchmark::benchmark)
5252
target_link_libraries(matrix_multiplication.cpp.x -ltbb)
5353
add_test(NAME matrix_multiplication.cpp.x COMMAND matrix_multiplication.cpp.x)
5454

55+
add_executable(algorithms_fft.cpp.x algorithms_fft.cpp)
56+
target_link_libraries(algorithms_fft.cpp.x fftw3)
57+
target_include_directories(algorithms_fft.cpp.x PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../include)
58+
add_test(NAME algorithms_fft.cpp.x COMMAND algorithms_fft.cpp.x)
59+
5560
add_executable(algorithms_gemm.cpp.x algorithms_gemm.cpp)
5661
target_link_libraries(algorithms_gemm.cpp.x benchmark::benchmark Boost::unit_test_framework)
5762
target_link_libraries(algorithms_gemm.cpp.x -ltbb)

benchmark/algorithms_fft.cpp

Lines changed: 127 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,127 @@
1+
// Copyright 2024-2026 Alfredo A. Correa
2+
// Distributed under the Boost Software License, Version 1.0.
3+
// https://www.boost.org/LICENSE_1_0.txt
4+
5+
#ifdef COMPILATION_INSTRUCTIONS
6+
g++ - std = c++ 17 - O3 - march = native - DNDEBUG - I../ include $0 - o $0.x - lfftw3 &&./ $0.x && gnuplot algorithms_fft_plots.gp;
7+
exit
8+
#endif
9+
10+
// Size sweep of multi::fft_plan vs FFTW, emitting gnuplot-friendly .dat files
11+
// (one per dimensionality). Methodology: both libraries build their plan once
12+
// (untimed) and recycle it; only execution is timed; CPU cache flushed before
13+
// every repetition; FFTW_ESTIMATE, no wisdom; single thread.
14+
#include <boost/multi/algorithms/fft.hpp>
15+
#include <boost/multi/array.hpp>
16+
17+
#include <fftw3.h>
18+
19+
#include <algorithm>
20+
#include <chrono>
21+
#include <cmath>
22+
#include <complex>
23+
#include <cstdio>
24+
#include <random>
25+
#include <vector>
26+
27+
namespace multi = boost::multi;
28+
using complex = std::complex<double>;
29+
30+
struct watch {
31+
std::chrono::high_resolution_clock::time_point s = std::chrono::high_resolution_clock::now();
32+
auto sec() const { return std::chrono::duration<double>(std::chrono::high_resolution_clock::now() - s).count(); }
33+
};
34+
static std::vector<char> g_thrash(64 << 20);
35+
void flush_cache() {
36+
for(std::size_t i = 0; i < g_thrash.size(); i += 64) {
37+
g_thrash[i]++;
38+
}
39+
char volatile x = g_thrash[0];
40+
(void)x;
41+
}
42+
template<class F> double time_it(long reps, F f) {
43+
double t = 0;
44+
for(long r = 0; r != reps; ++r) {
45+
flush_cache();
46+
watch w;
47+
f();
48+
t += w.sec();
49+
}
50+
return t / reps;
51+
}
52+
53+
auto reps_for(double n_total) -> long {
54+
double const work = 5.0 * n_total * std::log2(std::max(n_total, 2.0));
55+
return std::clamp<long>(static_cast<long>(2e8 / work), 5, 300);
56+
}
57+
58+
template<std::ptrdiff_t D>
59+
void sweep(std::vector<int> const& sides, char const* fname, char const* label) {
60+
std::FILE* out = std::fopen(fname, "w");
61+
std::fprintf(out, "# %s: multi::fft_plan vs FFTW 3 (plan recycled, exec only, cache flushed/rep, single thread, FFTW_ESTIMATE)\n", label);
62+
std::fprintf(out, "# mflops = 5*N*log2(N)/time_us (benchFFT convention), N = total points\n");
63+
std::fprintf(out, "# n_side N_total mine_ms fftw_ms mine_mflops fftw_mflops ratio_mine_over_fftw\n");
64+
for(int n : sides) {
65+
long N = 1;
66+
for(int d = 0; d != D; ++d) {
67+
N *= n;
68+
}
69+
std::vector<complex> base(N);
70+
std::mt19937 gen(42);
71+
std::uniform_real_distribution<double> dist(-1.0, 1.0);
72+
for(auto& e : base) {
73+
e = complex{dist(gen), dist(gen)};
74+
}
75+
long const reps = reps_for(static_cast<double>(N));
76+
77+
multi::array<complex, 1> flat(multi::extensions_t<1>{N});
78+
auto load = [&] { std::copy(base.begin(), base.end(), flat.begin()); };
79+
double mine = 0.0;
80+
{
81+
std::array<multi::ssize_t, D> ext_arr{};
82+
ext_arr.fill(n);
83+
multi::fft_plan<complex, D> const plan{ext_arr, multi::fft_forward};
84+
if constexpr(D == 1) {
85+
load();
86+
plan(flat);
87+
mine = time_it(reps, [&] { load(); plan(flat); });
88+
} else if constexpr(D == 2) {
89+
multi::array_ref<complex, 2> v(flat.data_elements(), {n, n});
90+
load();
91+
plan(v);
92+
mine = time_it(reps, [&] { load(); plan(v); });
93+
} else {
94+
multi::array_ref<complex, 3> v(flat.data_elements(), {n, n, n});
95+
load();
96+
plan(v);
97+
mine = time_it(reps, [&] { load(); plan(v); });
98+
}
99+
}
100+
101+
auto* in = static_cast<fftw_complex*>(fftw_malloc(sizeof(fftw_complex) * N));
102+
auto* fo = static_cast<fftw_complex*>(fftw_malloc(sizeof(fftw_complex) * N));
103+
auto loadf = [&] { for(long i = 0; i != N; ++i) { in[i][0] = base[i].real(); in[i][1] = base[i].imag(); } };
104+
fftw_forget_wisdom();
105+
fftw_plan p = D == 1 ? fftw_plan_dft_1d(n, in, fo, FFTW_FORWARD, FFTW_ESTIMATE)
106+
: D == 2 ? fftw_plan_dft_2d(n, n, in, fo, FFTW_FORWARD, FFTW_ESTIMATE)
107+
: fftw_plan_dft_3d(n, n, n, in, fo, FFTW_FORWARD, FFTW_ESTIMATE);
108+
loadf();
109+
double ffw = time_it(reps, [&] { loadf(); fftw_execute(p); });
110+
fftw_destroy_plan(p);
111+
fftw_free(in);
112+
fftw_free(fo);
113+
114+
double const work = 5.0 * static_cast<double>(N) * std::log2(static_cast<double>(N));
115+
std::fprintf(out, "%8d %10ld %12.5f %12.5f %12.1f %12.1f %8.3f\n", n, N, mine * 1e3, ffw * 1e3, work / (mine * 1e6), work / (ffw * 1e6), mine / ffw);
116+
std::fflush(out);
117+
std::fprintf(stderr, "%s n=%d done (reps=%ld)\n", label, n, reps);
118+
}
119+
std::fclose(out);
120+
}
121+
122+
int main() {
123+
sweep<1>({16, 32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384, 32768, 65536, 131072, 262144, 524288, 1048576, 2097152}, "fft_bench_1d.dat", "1D n");
124+
sweep<2>({16, 32, 64, 128, 256, 512, 1024, 2048}, "fft_bench_2d.dat", "2D n x n");
125+
sweep<3>({8, 16, 32, 64, 128, 256}, "fft_bench_3d.dat", "3D n x n x n");
126+
return 0;
127+
}

benchmark/algorithms_fft_plots.gp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
# Plots multi::fft_plan vs FFTW 3 from fft_bench_{1d,2d,3d}.dat
2+
# usage: run algorithms_fft.cpp.x first (writes the .dat files to the
3+
# current directory), then: gnuplot algorithms_fft_plots.gp
4+
# data columns: n_side N_total mine_ms fftw_ms mine_mflops fftw_mflops ratio
5+
6+
set terminal pngcairo size 900,600 font ",11"
7+
set grid xtics ytics mytics
8+
set logscale x 2
9+
set format x "2^{%L}" # x tics as powers of two
10+
set xlabel "transform size n (per axis)"
11+
set ylabel "speed 5·N·log_2(N) / time [MFLOPS] (higher is better)"
12+
set key top left
13+
set style line 1 lw 2 pt 7 ps 1.1 lc rgb "#c0392b" # multi
14+
set style line 2 lw 2 pt 5 ps 1.1 lc rgb "#2c3e50" # fftw
15+
16+
set format x "%.0f"
17+
set xtics rotate by -30
18+
19+
set output "fft_bench_1d.png"
20+
set title "1-D complex FFT — multi::fft\\_plan vs FFTW 3 (plan recycled, single thread)"
21+
plot "fft_bench_1d.dat" using 1:5 with linespoints ls 1 title "multi::fft\\_plan", \
22+
"fft_bench_1d.dat" using 1:6 with linespoints ls 2 title "FFTW 3 (estimate)"
23+
24+
set output "fft_bench_2d.png"
25+
set title "2-D complex FFT (n × n) — multi::fft\\_plan vs FFTW 3"
26+
plot "fft_bench_2d.dat" using 1:5 with linespoints ls 1 title "multi::fft\\_plan", \
27+
"fft_bench_2d.dat" using 1:6 with linespoints ls 2 title "FFTW 3 (estimate)"
28+
29+
set output "fft_bench_3d.png"
30+
set title "3-D complex FFT (n × n × n) — multi::fft\\_plan vs FFTW 3"
31+
plot "fft_bench_3d.dat" using 1:5 with linespoints ls 1 title "multi::fft\\_plan", \
32+
"fft_bench_3d.dat" using 1:6 with linespoints ls 2 title "FFTW 3 (estimate)"

include/boost/multi/algorithms/fft.NOTES.md

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -396,3 +396,41 @@ expressing butterflies through strided iterators would block the tight,
396396
auto-vectorizable batched loops that make this competitive with FFTW. The
397397
boundary between the two is the gather/scatter (or, on the fast paths, a raw
398398
`base()` pointer handed across) — FFTW draws the same line.
399+
400+
---
401+
402+
## 8. CUDA adaptation plan (not yet implemented)
403+
404+
The architecture is already CUDA-shaped: Stockham (no bit-reversal, uniform
405+
stages) is what GPU FFT libraries use; the batched layout `buf[k*m + j]`
406+
(batch contiguous) is exactly warp-coalesced when `j` maps to `threadIdx.x`;
407+
the fused strided stages' (pointer, stride, width) interface is
408+
layout-agnostic; and plans already separate "build tables once" (host) from
409+
"execute many" (upload once, launch many). Incremental steps:
410+
411+
1. **Extract butterfly bodies** into `BOOST_MULTI_HD` inline functions (the
412+
per-(block, r, j) work) shared by the host loops and device kernels. Pure
413+
refactor; CPU codegen must not regress (re-benchmark).
414+
2. **`fft_memory_space<Ptr>` trait** dispatching host (raw `T*`, current
415+
paths) vs device (thrust/Multi-CUDA fancy pointers), specialized in an
416+
opt-in adaptor header so the core stays CUDA-free. The current
417+
host-iterator gather fallback must *never* run on device pointers.
418+
3. **Device table mirror**: keep host `std::vector` tables; add a lazily
419+
uploaded POD `engine_view` (raw device pointers + sizes) per engine.
420+
4. **Device stage launchers**: one `__global__` kernel per stage kind around
421+
the shared butterfly body; thread → (j, r, block); stream-ordered launches
422+
replace the host stage loop. Grids are natively 3-D, so an extra
423+
(outer-slab, stride) index runs a whole 3-D middle axis in one launch.
424+
5. **Granularity flip**: same rotation/rank-descent orchestration, but the
425+
device branch batches the whole axis (m = all fibers; `mb_` is a host
426+
cache concern) and skips slab pairing (a CPU cache optimization). Needs a
427+
cap-and-tile fallback for the 2·n·m device scratch on huge arrays.
428+
6. **Bluestein**: pointwise chirp/kernel multiplies as elementwise kernels;
429+
sub-engine recursion stays host-side (it only orchestrates launches).
430+
**Six-step**: skip on device (its cache-blocking rationale is CPU-only).
431+
7. **Tests**: mirror `algorithms_fft.cpp` on Multi-thrust device arrays (the
432+
test already compiles as CUDA in the nvcc lane); `thrust::complex` should
433+
work unchanged through the `T{re,im}`/`+`/`-`/`*` contract and `fft_ops`.
434+
8. **Benchmark vs cuFFT** with the `adaptors/cufft` harness — the goal is
435+
*generic* device FFT (custom element types, arbitrary strided layouts,
436+
header-only), not beating cuFFT's tuned codelets for `float/double`.

0 commit comments

Comments
 (0)