SYCL Mode (Intel GPU)

GPU Mode 시리즈는 이름부터 CUDA mode 로 출발했다. 그런데 25강을 넘어서면서 — Mark 가 이름을 굳이 GPU Mode 로 바꾼 다음 첫 외부 백엔드 강의로 — Intel 의 Patric Zhao 를 불렀다. 강의의 이름이 일부러 “SYCL mode” 다.

강의가 답하려는 질문은 두 줄로 압축된다.

1. NVIDIA 가 아닌 GPU 위에서 같은 종류의 커널을 짜려면 어떤 모델이 있는가 — 그리고 그 모델이 CUDA 와 정확히 어디가 같고 어디가 다른가.
2. 그 모델로 LLM inference 같은 실전 워크로드를 어디까지 짜냈는가 — Intel GPU 위에서 fused MHA 와 flash attention 을 어떻게 구현했는지.

그래서 강의의 구조도 정확히 그 두 축이다. 전반부는 SYCL 의 이종 컴퓨팅 모델(host + device, queue, accessor, parallel_for) 을 CUDA 의 친숙한 모델 위에 mapping 한다. 후반부는 Intel oneAPI 위에서 SYCL 로 짠 fused operator 들 — RMSNorm + matmul, flash attention, MHA fusion — 의 실측 결과를 보여준다.

그래서 강의의 끝에 손에 잡혀 있어야 할 것은 — SYCL 의 핵심 개념 4개(queue, parallel_for/handler, USM, accessor), Intel GPU 의 메모리/스레드 모델(EU, slice, SLM), 그리고 이 모델이 CUDA 와 어디서 다른가 의 한 페이지짜리 mapping 이다.

강의의 첫 30분은 의도적으로 “heterogeneous computing” 이라는 더 넓은 단어로 시작한다. host + device 의 분리된 메모리, latency hiding 을 위한 SIMT 모델, async dispatch — 이 다섯 가지는 아키텍처 선택과 무관하게 모든 가속기에 공통이다.

이 frame 위에서 SYCL 이 등장하는 위치가 자연스러워진다 — 이 mapping 자체를 한 번에 묶는 single-source C++ layer. 같은 parallel_for 가 NVIDIA · Intel · AMD · CPU 위에서 backend 만 바꿔 돈다는 약속.

SYCL 의 spec 자체는 두꺼운 문서지만, 강의에서 Patric 이 잡은 진입점은 네 개의 추상이다. 이 넷만 알면 vector add 부터 GEMM 까지 다 짤 수 있다.

// SYCL vector add — 강의의 첫 코드
#include <sycl/sycl.hpp>
using namespace sycl;

int main(){
  constexpr int N = 1024;
  queue q{gpu_selector_v};

  // 1. USM 으로 device 메모리 할당
  float* a = malloc_shared<float>(N, q);
  float* b = malloc_shared<float>(N, q);
  float* c = malloc_shared<float>(N, q);

  // 2. submit + parallel_for
  q.submit([&](handler& cgh){
    cgh.parallel_for(range<1>{N}, [=](id<1> i){
      c[i] = a[i] + b[i];
    });
  }).wait();

  free(a, q); free(b, q); free(c, q);
}

SYCL 자체는 vendor 중립이지만, Intel GPU(Arc, Data Center GPU Max, Ponte Vecchio) 위에서는 Xe architecture 의 단어들을 다시 만난다. 이 자리에 익숙해져야 SYCL 커널의 launch 설정이 직관적으로 잡힌다.

// sub-group 사용 — Intel GPU 의 native warp 폭 활용
q.submit([&](handler& cgh){
  cgh.parallel_for(
    nd_range<1>{global, local},
    [=](nd_item<1> it)
      [[intel::reqd_sub_group_size(16)]] {
        auto sg = it.get_sub_group();
        int id  = sg.get_local_id();
        // sub-group 내 reduce — CUDA shfl_xor 동치
        float v   = data[it.get_global_id(0)];
        float sum = reduce_over_group(sg, v, plus<>{});
        if (id == 0) atomic_ref(out[0])
          .fetch_add(sum);
      });
});

SYCL 은 Khronos 표준이고, DPC++ (Data Parallel C++) 는 Intel 이 그 표준을 구현한 컴파일러 + 확장이다. oneAPI 는 DPC++ 위에 쌓인 vendor 라이브러리 묶음 — 정확히 CUDA 위의 cuBLAS / cuDNN / NCCL 자리다.

2024 년 후반부터 PyTorch 본체에 XPU backend 가 합류했다. CUDA 의 torch.cuda 처럼 torch.xpu.is_available() 가 동작한다. 그 이전 길은 IPEX (Intel Extension for PyTorch) — 같은 작업을 plug-in 형태로 했다.

# PyTorch XPU backend (≥ 2.4)
import torch

assert torch.xpu.is_available()
device = "xpu:0"

x = torch.randn(8192, 8192, device=device)
y = torch.randn(8192, 8192, device=device)
z = x @ y                       # oneDNN GEMM 호출됨
torch.xpu.synchronize()

# IPEX 로 inference 가속 — fused op 자동 매칭
import intel_extension_for_pytorch as ipex
model = ipex.optimize(model, dtype=torch.float16)
out   = model(x.to("xpu"))

강의의 큰 지렛대 중 하나는 — Patric 이 청중을 “CUDA 알고 있다” 로 가정한다는 것. 그래서 곳곳에서 CUDA → SYCL 직역 을 던진다. 한 페이지로 정리해두면 코드 읽기가 즉시 빨라진다.

이 mapping 위에서 한 줄로 더 짚어야 하는 점 — SYCL 은 dependency 를 queue 가 자동 추적한다. CUDA 에서 stream 위 같은 work 는 자동 직렬화되지만, 다른 stream 으로 분리하면 이벤트로 동기화를 사용자가 한다. SYCL 은 같은 buffer/USM 에 대한 read/write 의존을 보고 알아서 묶는다 — “잘못된 동기화로 race 가 나는 일이 줄어든다” 가 강의의 작은 메시지.

강의 후반부의 본론 — Intel GPU 위에서 LLM inference 의 hot path 를 SYCL 로 다시 짠 사례. fused RMSNorm + matmul, flash attention v2, MHA fusion. CUDA 쪽의 같은 개념과 같은 자리에 박힌다.

SYCL 의 가장 큰 마케팅은 vendor portability다. Intel · NVIDIA · AMD · 일부 NPU. 그런데 “portable 하지만 portable performance 는 보장 안된다” 는 갭이 늘 따라온다. 강의에서 Patric 이 이 갭을 명시적으로 인정한다.

강의에서 6개월 뒤 돌아왔을 때 가장 빨리 복원해야 하는 사실들과 — 직접 손에 박아야 하는 코드 자료들.

손에 새기기 — 실습 시퀀스

1. SYCL vector add 를 한 번 빌드 — Intel oneAPI 또는 Codeplay CUDA plugin 으로. icpx -fsycl vector_add.cpp -o vec. 같은 코드를 NVIDIA backend 로 빌드해서 두 device 위에서 같은 결과를 얻는다.
2. CUDA → SYCL 직역 — 자기가 가지고 있는 CUDA 커널 한 개를 골라 SYCL 로 옮긴다. __shared__ → local_accessor, __syncthreads → barrier mapping 을 직접 손에 박는다.
3. PyTorch XPU 베이스라인 — Arc GPU 또는 cloud Intel GPU 위에서 torch.xpu 로 LLaMA-2 7B inference 를 돌린다. fp16 / bf16 으로 토큰당 latency 측정.
4. IPEX optimize 적용 — 같은 모델에 ipex.optimize 한 줄. fused op 가 켜졌을 때의 차이를 본다.
5. SLM 사용한 reduction — local_accessor<float> 로 work-group 내 reduction 을 직접 짠다. reduce_over_group 빌트인과 손으로 짠 버전의 perf 비교.
6. sub-group shuffle — XOR 패턴으로 warp-level reduction 의 SYCL 변형을 짠다. permute_group_by_xor.
7. flash attention 의 SYCL 버전 살펴보기 — oneDNN 또는 IPEX 의 attention 구현을 git clone 해서 SLM 사용 패턴, sub-group 폭 결정, K/V tile 크기를 직접 읽는다.

SYCL 자체는 GPU Mode 시리즈에서 외로운 강의지만, 그 안의 개념들 — fused kernel, flash attention, vendor portability — 은 시리즈 전체와 연결된다.

학습 노트로 정리하면서 비워둔 자리들 — 강의 안에서 부분적으로만 등장한 주제, 또는 후속 검증이 필요한 주제.

• SYCL 의 NVIDIA backend 가 만드는 PTX 와 nvcc 가 만드는 PTX 의 차이 — 같은 vector add 를 두 경로로 빌드해 PTX 를 diff. “DPC++ 가 LLVM 을 거치니 결국 같은 자리에 도착한다” 는 강의의 주장이 PTX 수준에서 확인되는지.
• XMX vs Tensor Core 의 throughput 등가성 — joint_matrix API 로 짠 GEMM 이 cuBLAS GEMM 의 몇 % 를 따 내는지. 강의에서는 oneDNN 차원에서만 측정.
• CUDA graph ↔ SYCL command_graph — launch overhead 가 dominant 한 decode 루프에서 SYCL graph 가 같은 효과를 내는지. 강의 Q&A 에서 “sycl graph 가 있다” 정도로만 언급됨.
• Triton 의 Intel GPU backend — Triton 3.x 시점에서 Intel GPU 가 정식 backend 로 들어오고 있다. 같은 LLM 커널을 Triton + XPU 로 짜면 SYCL 직접 작성 대비 어디까지 따라오는지.
• NPU / Habana 와의 관계 — 강의에서 “SYCL 이 NPU 도 노린다” 고 짧게 언급. 실제로 Gaudi 위에서 SYCL 코드가 어디까지 도는지 확인 필요.
• PyTorch 의 XPU backend 가 CUDA 와 feature parity 를 언제 잡는지 — torch.compile / FSDP / torch.distributed 위 XPU 지원의 현재 상태와 로드맵.