Data Processing on GPUs

GPU Mode 시리즈의 다른 강의들이 거의 다 모델 학습/추론 에 초점을 맞췄다면, 이 강의는 명시적으로 “데이터 엔지니어링이 GPU 에서 어떻게 일어나는가” 를 깐다. SQL 쿼리, ETL 파이프라인, dataframe API — pandas 같은 인터페이스인데 backend 는 CUDA.

강의가 깐 큰 질문 세 개.

1. 왜 데이터 처리를 GPU 에서 하는가 — 의외로 “GPU 가 빠르니까” 가 아니다. 정확한 답은 HBM 의 대역폭(§ 02).
2. CSV 같은 가변길이 레코드를 GPU 에서 어떻게 파싱하는가 — sequential 하게 line-by-line 도는 게 inherent 한 알고리즘인데, GPU 는 sequential 을 못 한다. scan 으로 푼다(§ 04).
3. groupby/aggregation 같은 SQL 기본기는 어떻게 짜는가 — atomic + GPU hashmap 의 조합(§ 05–06).

강의 도입의 한 figure. CPU DRAM 의 대역폭과 GPU HBM 의 대역폭을 같은 축에 놓는다.

그런데 데이터가 처음부터 GPU 위에 있을 리는 없다 — disk → CPU memory → GPU memory 의 PCIe 단계를 지난다. 그래서 “disk 에서 한 번만 읽고, 그 다음 모든 처리는 GPU 안에서” 의 패턴이 표준. cuDF 의 read_csv · read_parquet 가 처음부터 GPU 메모리로 직접 적재하는 이유다.

행렬 곱 같은 GPU 일은 — “element 위치가 자기 thread 가 어디 있는지” 가 곧바로 결정된다. CSV 는 다르다. line N 이 어디서 시작하는지를 알려면 line 0..N-1 을 다 읽어야 한다. inherent 한 sequential dependency.

강의의 작은 트릭 — bytestream 을 청크로 나누고 각 청크의 시작 부근에서 가까운 \n 까지 skip 한 자리를 row 시작으로 본다. partial line 은 두 청크의 경계에서 합치기. 이렇게 하면 대략적인 parallelism 은 얻지만 정확한 row index 는 아직 모른다 — 거기서 scan 이 들어온다.

강의에서 Devavret 이 가장 흥미롭다고 강조한 부분. 동료의 PhD thesis 였던 아이디어 — “CSV/JSON 파싱 전체를 scan 의 합성으로 표현”. 핵심은 한 byte 의 “role” (newline, separator, quote, escape, …) 이 inputs 자체에서만 결정되고, 그 role 의 누적이 곧 row/col index 라는 사실.

quote 안에 들어 있는 comma 는 어떻게 처리하나? — quote 의 toggle 도 scan 으로 표현된다 (XOR 누적). 강의는 깊이 들어가지 않지만, 같은 idea 의 확장.

SQL 의 가장 흔한 연산 — SELECT city, AVG(temp) FROM measurements GROUP BY city. 이걸 GPU 에서 짜는 표준 패턴이 hash → atomic add. 강의의 핵심 두 figure 중 하나.

강의의 한 가지 디테일 — bucket 수가 너무 작으면 collision 이 폭발하고, 너무 크면 메모리 낭비. 일반적인 cuDF/cuco 의 default 는 load factor 50% 정도로 bucket array allocate. dictionary 처럼 unique key 를 미리 알면 perfect hash 도 가능.

강의에서 Devavret 이 강조한 메시지. CPU 의 std::unordered_map 같은 chained hashmap 은 GPU 에서 잘 안 도는다 — pointer chasing 이 random read 를 만든다. GPU 는 항상 open addressing 의 flat array.

// GPU hashmap insert — 핵심만
__device__ void insert(slot_t* table, int capacity, K key, V val) {
  int idx = hash(key) % capacity;
  while (true) {
    // CAS — empty slot 을 잡으면 내가 차지
    K prev = atomicCAS(&table[idx].key, EMPTY, key);
    if (prev == EMPTY) {       // 새 entry — value 초기화
      atomicAdd(&table[idx].value, val);
      return;
    }
    if (prev == key) {         // 이미 같은 key 의 entry — 누적
      atomicAdd(&table[idx].value, val);
      return;
    }
    idx = (idx + 1) % capacity; // linear probe
  }
}

그리고 강의의 또 한 자리 — shared memory 안 hashmap. 작은 cardinality (수백~수천 unique) 면 block 하나가 자기 hashmap 을 SRAM 안에 두고 끝낸다. 그 다음 block 들의 결과를 hierarchy 로 합친다. cuDF 의 groupby 가 정확히 이 패턴.

강의의 가장 구체적인 케이스 스터디. 1 Billion Row Challenge — 10억 줄의 (도시, 온도) 측정값에서 도시별 min/avg/max 를 뽑아라. CPU 에서 ~10초 안에 푸는 게 목표. GPU 에서 두 접근.

강의에서 보여준 결과 — tailor-made 가 더 빠르긴 하지만 차이가 생각만큼 크진 않다. SQL engine 도 well-engineered 면 데이터 통과를 1–2 회로 압축한다. “손으로 짠 커널의 이득은 알고리즘적 통합이고, SQL 엔진의 이득은 generality” 의 trade-off 가 명확히 드러나는 사례.

CSV 가 가변길이의 row-oriented 면, Parquet 는 columnar + 사전 디코드 가능한 메타데이터를 들고 있다. row group / page / dictionary 의 계층. 각 column 이 contiguous 하게 저장되어 있어 GPU coalesced load 와 자연스럽게 맞는다.

그리고 dictionary encoding — column 의 값이 적은 unique 면, 사전(dictionary) 한 번 + 인덱스 array. 압축 + lookup 효율 둘 다 좋아진다. 강의에서 Devavret 이 “cuDF Parquet writer 의 dictionary encoding 커널을 다시 짰던” 일화 (§ 06 참고) 가 이 자리.

Devavret 의 현재 회사 Voltron Data 의 제품 — Theseus / Velox 같은 GPU-native query engine. SQL 을 받아서 GPU 친화적 execution plan 으로 컴파일해서 돌린다. cuDF 가 “손으로 호출하는 라이브러리” 라면 Voltron 은 “DB 처럼 쓴다”.

이 영역의 큰 통찰 — “데이터 엔지니어링이 GPU 로 옮겨가는 건, ML 이 GPU 로 옮겨간 패턴의 5–10년 뒤 반복”. 같은 hardware/software 스택을 다른 도메인에서 다시 쓰는 흐름.

데이터 엔지니어링이 GPU 에서 어떻게 돌아가는지의 첫 mental model — 그 핵심 사실들.

손에 새기기 — 실습 시퀀스

1. cuDF first run — 1 GB 정도의 CSV 를 pandas 와 cuDF 로 같은 groupby 돌려보고 시간 비교. 차이가 30× 가까운지 확인.
2. scan 으로 \n 위치 찾기 — 작은 byte stream 위에서 newline 위치를 직접 prefix-sum 으로 누적하는 CUDA 커널 짜기. cub::DeviceScan 활용.
3. shared memory hashmap — 100 unique key 에 대해 block 하나가 SRAM 위에 hashmap 하나 두고 atomic 누적. 이후 block 간 합치기.
4. 1BRC tailor kernel — 자기가 직접 작은 1BRC clone (10M rows) 을 짜고 — SQL/pandas vs cuDF vs tailor kernel 세 시간 비교.
5. cuCollections 직접 써보기 — cuco::static_map 으로 unique counting 같은 작업. open addressing 의 insert 함수 source 직접 읽기.
6. Parquet page-level parallel — 작은 Parquet 파일의 column chunk 안 page 수를 footer 에서 읽고, page 단위 parallel decode 가 가능한지 head 만 본다.

• scan-based CSV parser 의 정확한 구현 — 강의에서 reference 한 PhD thesis/paper 의 정확한 출처는 자막에서 명시되지 않음. 추적 필요 (확인 필요).
• cuDF 의 groupby kernel 위치 — repo 의 cpp/src/groupby/hash/ 안 코드를 직접 읽으면 § 05–06 의 mental model 과 정확히 어떻게 매칭되는지 확인 가능.
• 1BRC 의 GPU 결과 — 강의에서 정확한 시간 숫자가 명시되지 않음. tailor-made vs cuDF 의 차이 % 가 얼마인지 자기 GPU 에서 재현 필요.
• Voltron 의 internals — Theseus 가 actually 어떻게 SQL → GPU plan 을 만드는지 강의에서 깊이 들어가지 않음. 별도로 paper / blog 추적.
• JSON 파싱의 stack-based 분해 — § 04 의 callout 에서 언급한 nesting depth scan. simdjson 의 알고리즘과 GPU 버전의 차이는 후속 자료로 보강 필요.
• multi-GPU groupby — 강의는 single-GPU 위주. multi-GPU 에서 hash partition + shuffle 이 어떻게 도는지는 RAPIDS 의 dask-cuDF 영역.