Apple 실리콘 GPU로 vDSP 제압, 138 GFLOPS 래디스‑8 FFT 구현

본 논문은 Apple Silicon GPU, 특히 M1 칩에 최적화된 Fast Fourier Transform(FFT) 구현을 제시한다. 연구 동기는 레이더·SAR 등에서 대규모 1차원 FFT가 빈번히 사용되며, 기존 Apple의 vDSP/Accelerate 라이브러리가 GPU 가속을 충분히 활용하지 못한다는 점이다. 이를 해결하기 위해 저자들은 2015년 Intel 통합 GPU용 FFT 최적화 논문(두 단계 로컬 메모리 모델)을 참고해 Apple GPU의 고유 구조에 맞는 ‘두 단계 로컬 메모리 모델’을 정의한다. 첫 번째 단계는 208 KiB 레지스터 파일을 데이터가 상주하는 영역으로 활용하고, 두 번째 단계는 32 KiB 스레드그룹 메모리를 교환 전용 파이프라인으로 제한한다. 이 모델은 레지스터와 SIMD shuffle을 최대한 활용해 메모리 접근을 최소화하고, 스레드그룹 배리어 비용이 매우 낮아(≈2 사이클) 배리어 최소화보다 접근 패턴 최적화가 핵심임을 보여준다. 핵심 알고리즘은 래디스‑8 split‑radix Decimation‑In‑Time(DIT) Stockham 구조이다. N=4096, 배치 256을 기준으로 512 스레드(16 SIMD 그룹)로 구성된 커널은 4단계만에 변환을 완료한다. 각 스레드는 8개의 복소 샘플을 레지스터에 보관하고, 52 실수 덧셈·12 실수 곱셈으로 구성된 효율적인 8‑점 DFT 버터플리를 수행한다. 트위들 인자는 단일 sincos 호출 후 복소 곱셈을 재귀적으로 적용해 계산량을 크게 줄였다. 이 설계는 레지스터 사용량을 약 30 % 수준으로 유지하면서도 138.45 GFLOPS(1.78 µs/FFT)를 달성, vDSP의 107 GFLOPS 대비 29 % 향상을 기록한다. 또한, 래디스‑4 Stockham 커널(6패스, 1024 스레드)과 SIMD‑shuffle 기반 변형(래디스‑32)도 구현했으며, 각각 113.6 GFLOPS와 61.5 GFLOPS를 보였다. 특히 SIMD‑shuffle 변형은 스레드그룹 메모리의 스트라이드 접근이 3배 느리다는 실험 결과와 맞물려 성능이 크게 저하된 것을 확인했다. 이는 두 단계 모델이 ‘접근 패턴 최소화’를 우선시해야 함을 실증한다. Apple 전용 하드웨어 매트릭스 연산인 simdgroup_matrix(8×8 MMA)도 탐색했지만, 데이터 마샬링 오버헤드가 커서 단일 FFT당은 실질적인 이득이 없었다. 다만 배치가 큰 경우(예: 256~16384 FFT)에는 여러 FFT를 동시에 MMA에 매핑해 효율을 높일 가능성이 제시되었다. 크기 확장 측면에서는 레지스터와 스레드그룹 메모리 용량을 고려해 한 스레드그룹당 최대 4096 포인트 FFT를 직접 처리하도록 설계했다. N이 4096을 초과하는 경우에는 4‑step FFT(분할‑정복) 방식을 적용해 N=N₁·N₂ (N₂≤4096) 형태로 두 단계 디스패치를 수행하고, 중간 단계에서 디바이스 메모리 전이를 이용해 트위들을 적용한다. Apple의 통합 메모리 모델 덕분에 전이 비용이 기존 이산 GPU 대비 크게 감소한다. 실험 결과, N=8192와 N=16384에 대해 각각 112 GFLOPS와 103 GFLOPS를 기록했으며, 이는 여전히 vDSP보다 높은 수준이다. 배치 규모에 따른 스케일링 실험에서는 배치가 64 이상일 때 GPU 커널이 GPU 코어를 충분히 활용해 vDSP를 앞선다. 배치가 작을 경우(vDSP는 디스패치 오버헤드가 낮아) vDSP가 우위를 점한다. 결론적으로, 이 논문은 Apple Silicon GPU의 레지스터‑우선 로컬 메모리 구조를 정량화하고, 이를 기반으로 설계된 래디스‑8 Stockham FFT가 기존 CPU 기반 vDSP보다 현저히 높은 실효 성능을 달성함을 입증한다. 향후 연구 과제로는 다중 FFT 배치를 이용한 MMA 활용, 더 큰 N에 대한 다단계 전이 최적화, 그리고 다른 Apple GPU(예: M2, M3)에서의 확장성을 검증하는 것이 제시된다.