gpumode · 강의 아카이브
《GPU Mode》 L081 2025 High priority transcript · failed

High-performance Purely Functional Data-Parallel Array Programming

Triton 과 ThunderKittens 의 풍경 옆에 전혀 다른 길 — Troels Henriksen 이 만든 Futhark. 사용자는 map/reduce/scan 같은 SOAC (Second-Order Array Combinator) 만 적고, AOT 컴파일러가 GPU 코드로 떨어뜨린다. CUDA 의 lifetime 문제 없이 nested parallelism 을 표현하는 functional 답안의 학습 노트.

Futhark SOAC functional data-parallel Copenhagen prefix sum flat parallelism array language
T
Speaker
Troels Henriksen
University of Copenhagen · Futhark 개발 lead
강의 번호
L081
학습 우선순위
High
자막
failed
출신
academic / open
§ 01강의가 풀려는 문제· why this lecture exists

"GPU 프로그래밍은 왜 매번 같은 boilerplate 인가" 의 다른 답

CUDA / Triton 의 풍경에서 사용자는 거의 항상 같은 일을 반복한다 — index 계산, tile 분할, shared memory 의 lifetime 관리, race 회피. 이 boilerplate 가 본질적인가, 아니면 잘못된 추상의 부산물인가?

Futhark 는 후자의 답에 베팅한다. functional 추상 (map, reduce, scan, segmented reduce) 으로 표현된 알고리즘은 컴파일러가 자동으로 효율적 병렬 코드로 lowering 할 수 있다. 사용자는 알고리즘만 적고 lifetime / race / vectorization 은 컴파일러의 책임.

강의의 frame.

  1. data-parallel 의 본질은 functional 이다 — map/reduce/scan 은 모두 closure-form 으로 정의 가능. side effect 가 없으면 compiler 가 reorder, fuse, parallelize 가능.
  2. nested parallelism 의 문제map (\x -> map f x) y 같은 nested 패턴이 자연스럽지만 GPU 의 grid/block 모델에 직접 mapping 안 됨. Futhark 의 flattening 이 이 자리를 푼다.
  3. AOT 컴파일 — Python jit 이 아닌 ahead-of-time. C / CUDA / OpenCL 코드 생성. 결과물이 일반 binary 처럼 deploy 가능.
강의의 인지적 frame

GPU 프로그래밍을 "자료구조 위에서 reduce/scan 의 합성" 으로 다시 본다. 이게 어색해 보일 수 있지만 — 사실 이미 익숙한 패러다임이다 (NumPy, JAX, MapReduce 가 모두 이 frame). Futhark 는 그것을 GPU 위에서 native 로 끌어내는 시도.

"GPU 의 모든 빠른 코드는 결국 reduce 와 scan 의 합성으로 쓸 수 있다 — 그러니 그걸 first-class 로 두자."Troels Henriksen · 확인 필요
§ 02functional array language 의 가치· why functional

side effect 가 없으면 compiler 가 자유롭다

Futhark 가 "purely functional" 임을 강조하는 이유 — side effect 가 없는 코드는 컴파일러가 거의 자유롭게 reorder, fuse, parallelize 할 수 있다.

왜 이게 GPU 에 특히 좋은가

GPU 는 본질적으로 SIMD/SIMT 모델 — 같은 일을 다른 데이터에 동시에. 이건 정확히 map 의 모양이다. functional 추상이 hardware 의 모양과 1:1. CUDA 의 thread index 계산은 그 1:1 mapping 을 사람이 매번 손으로 하는 것.

§ 03컴파일 모델· how Futhark lowers

C / CUDA / OpenCL 로 떨어지기까지의 단계

L0 Futhark source .fut 파일. ML-family 문법. SOAC 와 일반 함수의 합성
L1 SOACS IR map / reduce / scan 등 high-level 추상이 살아 있는 IR
L2 fusion + flattening nested parallelism 을 flat 으로 변환. 인접 SOAC 합치기
L3 low-level kernel IR flat kernel 들 — segmented reduce / map 의 unrolled 형태
L4 backend 선택 CUDA · OpenCL · multicore C · HIP · sequential C
L5 코드 생성 최종 .c / .cu 파일. nvcc / clang 으로 binary 빌드

이 pipeline 의 핵심 — fusion 과 flattening 이 가장 가치 있는 단계. fusion 은 중간 array 를 없애 memory bandwidth 절감, flattening 은 nested parallelism 을 GPU 의 grid/block 위에 mapping.

flattening 의 의미를 풀어본다.

  • regular nestedmap (\x -> map f x) xs — inner map 의 size 가 같으면 outer × inner 평면을 하나의 grid 로 펼친다. 자명.
  • irregular nested — inner size 가 다르면 (variable-sized rows) — segmented operations 으로 변환. row offset 배열을 들고 다닌다.
  • scan 의 nested — outer map 안의 scan. segmented scan 으로 변환.

이 단계의 어려움은 "flattening 이 항상 효율적이지 않다" 는 점. 작은 inner size 면 thread divergence 비용이 큼. 컴파일러가 heuristic 으로 결정.

-- Futhark source — sparse matmul
def sparse_matmul (rows: [][]i32) (vals: [][]f32) (x: []f32) =
  map2 (\row vs ->
    reduce (+) 0f32
    (map2 (\j v -> v * x[j]) row vs))
    rows vals

-- 컴파일러가 만드는 형태 (개념)
-- 1. nested map → segmented map
-- 2. inner reduce → segmented reduce
-- 3. flatten 으로 single kernel

-- 결과: 한 번의 kernel launch,
--       row 마다 다른 길이를 segment scan 으로 처리
§ 04GPU mapping· SOAC → kernel

map / reduce / scan 이 grid/block/thread 위로 떨어지는 방식

FIG · SOAC composition diagramfunctional → CUDA grid/block/thread
map (\x -> reduce (+) 0 (scan (+) 0 (map g x))) xs composed SOAC pipeline fusion fused: map (\x -> sum_of_prefix_sum_of_g x) xs single SOAC, inner = composed flatten + lowering grid outer map → blocks block per "x" size = N (outer) block inner scan → SMEM scan cooperative block-level reduce thread inner map (g) → register per-element work vectorized load
SOAC 의 composition 패턴이 GPU 의 3-level hierarchy 와 자연스럽게 만난다 — 외곽 map → grid, 중간 scan → block (cooperative), 안쪽 map → thread. 사용자는 이 mapping 을 직접 안 적는다.
map
map f xs
[a] → (a → b) → [b]
완전 독립 element-wise. 가장 자연스러운 GPU primitive.
reduce
reduce f e xs
[a] → (a → a → a) → a → a
associative 이면 tree-reduce 로 O(log n). f 는 monoid.
scan
scan f e xs
[a] → (a → a → a) → a → [a]
prefix sum 의 일반화. work-efficient parallel scan 구현.
filter
filter p xs
[a] → (a → bool) → [a]
scan 으로 표현 가능 — index 누적 후 compaction.
scatter
scatter dst is vs
[a] → [i32] → [a] → [a]
indexed write. uniqueness type 으로 in-place 가능.
stream_red
stream_red op f xs
chunk 단위 reduce
large array 의 sequential 한 chunk 처리 — register tiling.
§ 05prefix sum 등 빌딩블록· scan, reduce, segmented

"work-efficient parallel scan" 이 모든 것의 base

Futhark 의 모든 algorithm 이 결국 map/reduce/scan 의 합성. 그 중에서도 scan 이 핵심 — 한 단계 더 깊은 병렬화의 도구.

parallel scan 은 1980 년 Blelloch 의 work-efficient algorithm 으로 잘 알려져 있다. n 개 element 의 prefix sum 을 O(log n) depth 로, total work O(n) 으로 계산.

왜 scan 이 GPU 에 특히 좋은가

scan 은 local 작업 + tree reduction + broadcast 의 합성. CUDA 의 warp shuffle, block 단위 reduce 와 1:1. CUB, Thrust 같은 라이브러리가 모두 정교한 scan 구현을 가지고 있다 — Futhark 는 그것들을 사람이 매번 부르지 않게 추상.

"GPU 위에서 빠른 모든 알고리즘은 reduce 와 scan 의 합성이다 — Futhark 는 그것을 사용자의 손에서 뗀다."학습 노트
§ 06ML 적용· where Futhark fits

"Futhark 로 LLM 을 짤 수 있는가" 의 답

Futhark 의 ML 적용 가능성 — 강의에서 Troels 가 명시적으로 짚는 자리. 결론은 "가능하지만 지금은 PyTorch / JAX 만큼 큰 모델을 cover 하기는 어렵다".

그러나 한계도 명확하다.

"실용적 자리"

Futhark 가 빛나는 자리 — 커스텀 ML / 과학 계산 알고리즘. 예: 새 attention variant 의 reference, image processing pipeline, simulation. PyTorch 의 default op 안에 없는 자리에 prototype 으로 좋다.

§ 07채택 사례· production users

"academic 출신, 실제 production 에서 쓰이는가" 의 답

큰 그림 — Futhark 는 산업 default 가 아니지만 niche 의 strong tool. JAX/PyTorch 가 cover 안 하는 자리, 또는 functional 추상이 코드 정확성에 결정적인 자리.

§ 08다른 array language 와 비교· APL · Halide · JAX

같은 산을 오르는 다른 길들

language패러다임compileGPU사용처
Futharkpure functionalAOTCUDA / HIPHPC, custom ML, prototype
APL / J / Qarray calculusinterpreted제한적finance, data analysis
Halideschedule-separatedJIT/AOTCUDA / Metalimage processing
JAX (Pallas)functional + tracingJIT (XLA)TPU / GPUML research
Accelerate (Haskell)pure functionalJITCUDAacademic, niche
SaCsingle-assignment CAOT제한적academic
Tritontile DSLJITCUDAML kernel
NumPyimperative arrayinterpretedcuPy 별도data analysis, ML

핵심 비교 차원.

"functional array language 는 항상 niche 였다 — 대중화하지는 못했지만, 그 추상이 옳다는 사실은 JAX 의 존재로 증명됐다."학습 노트
§ 09한계· caveats

Futhark 가 못하는 것들

§ 10기억할 메모· key takeaways

다시 열었을 때 손에 잡혀야 할 것

SOAC
Second-Order Array Combinator. map/reduce/scan/scatter 등. Futhark 의 first-class primitive.
pure functional
side effect 없음. 컴파일러가 reorder, fuse, parallelize 자유. race condition 본질적 부재.
flattening
nested parallelism 을 grid/block 위 flat 으로. regular vs irregular 처리 다름.
fusion + deforestation
map (map f) → map (f . id). 중간 array 안 만든다. memory bandwidth 절감.
work-efficient scan
Blelloch up-sweep + down-sweep. O(n) work, O(log n) depth. 모든 알고리즘의 base.
uniqueness type
in-place update 의 functional escape. mutation 을 표현하면서도 race 안전.
AOT 컴파일
CUDA / OpenCL / HIP / multicore C 로 떨어짐. Python 모듈로도 export. binary deploy.
ML niche
새 op prototype, custom 과학 계산, research. cuDNN / FA4 같은 hand-tuned 자리는 못함.
YouTube강의 영상 (확인 필요)
RelatedHalide · Accelerate · SaC · Co-dfns · JAX

손에 새기기 — 실습 시퀀스

  1. hello-world — Futhark 설치, 1-D dot product. map2 (*) xs ys |> reduce (+) 0. CUDA backend 로 컴파일, Python 에서 호출.
  2. prefix sum 직접 짜기 — Blelloch scan 을 Futhark 의 scan (+) 0 으로. 같은 알고리즘을 raw CUDA 로 짜고 코드 길이 비교.
  3. fusion 효과 확인map f (map g xs) 의 IR dump. 컴파일러가 두 map 을 합쳤는지 직접 확인.
  4. segmented reduce — variable-row matrix 의 row-wise sum. flag array 로 segment 표시. CUB 의 SegmentedReduce 와 성능 비교.
  5. mini ML — 3-layer MLP 의 forward + backward 를 Futhark 로. autodiff vjp 사용. 결과를 PyTorch 와 atol 비교.
  6. raw CUDA backend 출력 읽기 — Futhark 의 --save-cuda 로 생성된 .cu 파일. 사람이 짠 CUDA 와 어떻게 다른지 직접.
§ 11다른 강의로의 연결· connections

이 강의가 시리즈 안에서 어디로 이어지는가

§ 12열린 질문· open questions

다음에 다시 들었을 때 직접 검증해야 할 것들

검증 메모

이 노트는 Futhark 공식 site / 도메인 지식 (functional array programming, parallel scan) / 강의 metadata 의 재구성. 모든 절대 수치는 시점 의존 — futhark-lang.org 와 책 직접 참조 권장. 강의의 정확한 demo 코드와 측정은 영상 직접 확인 필요.

← Lecture 080How FlashAttention 4 Works Lecture 082 →다음 강의로