필사 모드: AMD와 NVIDIA, 무엇이 다른가 — 하드웨어보다 스택이 문제인 이유

• 들어가며 — HIPIFY가 90퍼센트를 옮겨 준 뒤에 남는 10퍼센트
• 하드웨어 — SM과 CU, 텐서 코어와 매트릭스 코어
• 워프 32 대 웨이브프론트 64 — 가장 비싼 한 줄
• 소프트웨어 스택 — CUDA와 ROCm, HIP
  • 라이브러리 대응표
• 이식 경로 — HIPIFY가 되는 것과 안 되는 것
  • 자동으로 옮겨지는 것
  • 손으로 고쳐야 하는 것
  • 이식 후 반드시 해야 할 검증
• 현실적으로 무엇이 걸림돌인가
  • 1. 커널 생태계의 축적량
  • 2. 프레임워크 지원의 시차와 지원 매트릭스
  • 3. 문제가 생겼을 때의 정보량
• 어떤 워크로드라면 AMD가 합리적인가
• 마치며 — 격차는 실리콘이 아니라 축적된 커널에 있다
• 참고 자료

들어가며 — HIPIFY가 90퍼센트를 옮겨 준 뒤에 남는 10퍼센트

CUDA 코드베이스를 AMD로 옮기는 작업은 대체로 이렇게 흘러갑니다. hipify-perl을 돌리면 파일 대부분이 몇 초 만에 변환됩니다. cudaMalloc이 hipMalloc이 되고 __global__은 그대로이고 커널 실행 문법도 같습니다. 빌드가 통과하고, 작은 테스트가 통과하고, 여기까지 반나절입니다.

그다음 2주가 걸립니다. 결과가 미묘하게 틀리는 커널 하나를 찾다가 __shfl_xor 뒤에 32가 하드코딩되어 있는 것을 발견하고, 성능이 절반인 커널을 붙잡고 있다가 셰어드 메모리 타일 크기가 워프 32를 전제로 정해져 있었다는 것을 알게 됩니다. 인라인 PTX가 들어간 파일 하나는 통째로 다시 써야 합니다.

이 글은 그 10퍼센트가 정확히 무엇인지, 그리고 그것보다 훨씬 중요한 생태계의 차이가 무엇인지를 다룹니다. 응원도 비난도 없이 구조만 봅니다. 확인 기준은 ROCm 7.14.0(2026년 7월 16일 릴리스), vLLM v0.26.0 문서, AMD ROCm 공식 문서입니다.

하드웨어 — SM과 CU, 텐서 코어와 매트릭스 코어

용어부터 대응시킵니다. 개념 층위가 거의 그대로 맞습니다.

빌드할 때 대상 아키텍처를 명시해야 하는 것도 같습니다. NVIDIA에서 -arch=sm_90을 주듯 ROCm에서는 gfx942 같은 값을 줍니다. 자기 장비의 값은 rocminfo로 확인합니다.

rocminfo | grep gfx        # 예: gfx942 (MI300 계열), gfx950 (MI350 계열)
rocm-smi                   # nvidia-smi 대응. 사용률, 온도, 전력

지금 시장에서 마주치는 세대를 정리하면 이렇습니다. 아래 하드웨어 수치는 벤더 발표와 2차 보도를 종합한 것으로, 저는 실측하지 않았습니다. 구매 판단에는 반드시 벤더의 공식 스펙 시트를 확인하십시오.

여기서 읽어야 할 것은 개별 숫자가 아닙니다. 두 회사의 하드웨어가 같은 세대에서 대체로 비슷한 메모리 용량과 대역폭 급에 도달해 있다는 사실입니다. AMD가 용량에서 앞서는 세대도 있었고 NVIDIA가 따라잡은 세대도 있습니다. 이 축에서 결정적 격차는 없습니다.

그래서 실무의 격차는 다른 곳에서 옵니다. 그 다른 곳이 이 글의 나머지입니다.

워프 32 대 웨이브프론트 64 — 가장 비싼 한 줄

하드웨어 차이 중 코드에 직접 상처를 내는 것은 이 하나입니다. NVIDIA 워프는 32개 스레드입니다. AMD CDNA 계열 데이터센터 GPU의 웨이브프론트는 64개입니다. RDNA 계열은 32개입니다.

HIP 문서가 이 지점을 명시적으로 경고합니다.

Code should not assume a warp size of 32 or 64, as AMD GPU architectures have different warp sizes. The warpSize built-in should be used in device code.

이 한 줄이 실제로 만드는 버그는 세 종류입니다.

첫째, 셔플 리덕션의 반복 횟수. 앞선 글에서 쓴 워프 리덕션을 다시 보겠습니다.

// NVIDIA를 전제로 쓴 코드. AMD에서 절반만 리덕션된다.
__inline__ __device__ float warpSum(float v) {
  for (int off = 16; off > 0; off >>= 1)          // 32를 전제한 16
    v += __shfl_xor_sync(0xffffffffu, v, off);
  return v;
}

웨이브프론트가 64면 off가 32부터 시작해야 합니다. 16부터 시작하면 앞 32개 레인과 뒤 32개 레인이 각각 따로 합산되고, 최종 결과는 절반만 맞습니다. 컴파일도 되고 크래시도 안 나고 그냥 틀립니다. 이식 작업에서 가장 오래 붙잡게 되는 종류의 버그입니다.

이식 가능한 형태는 이렇습니다.

// HIP. AMD와 NVIDIA 양쪽에서 맞게 동작한다.
__device__ float waveSum(float v) {
  for (int off = warpSize / 2; off > 0; off >>= 1)
    v += __shfl_xor(v, off, warpSize);
  return v;
}

둘째, 레인 마스크의 폭. NVIDIA에서 활성 레인 마스크는 32비트라 0xffffffff가 자연스럽습니다. 웨이브프론트 64에서는 64비트가 필요합니다. HIP 문서는 여기에 구체적인 함정을 하나 더 적어 두었습니다. 64폭 웨이브프론트에서 32비트 정수를 31보다 큰 값만큼 시프트하면 레지스터가 0으로 지워진다는 것입니다. 해법은 레인 마스크에 uint64_t를 쓰는 것입니다.

셋째, 타일 크기와 블록 크기의 암묵적 전제. 블록 크기 256을 "8워프"로 생각하고 셰어드 메모리 부분합 배열을 8칸으로 잡아 두었다면, AMD에서는 웨이브프론트가 4개라 4칸만 쓰이고 로직이 어긋납니다. 반대로 32의 배수로만 맞춰 둔 블록 크기가 AMD에서는 64의 배수가 아니라 마지막 웨이브프론트의 절반이 놀 수 있습니다.

원칙은 하나입니다. 32도 64도 소스에 상수로 쓰지 않습니다. warpSize를 쓰거나, 정 필요하면 컴파일 타임 상수로 한 곳에 모아 두고 아키텍처별로 갈아 끼웁니다.

소프트웨어 스택 — CUDA와 ROCm, HIP

층을 나란히 놓으면 이렇습니다.

NVIDIA                          AMD
──────────────────────          ──────────────────────
PyTorch / JAX / vLLM            PyTorch / vLLM
   │                               │
cuBLAS, cuDNN, NCCL             rocBLAS, MIOpen, RCCL
   │                               │
CUDA Runtime API                HIP Runtime API
   │                               │
CUDA Driver                     ROCr 런타임 + ROCk 커널 드라이버
   │                               │
NVCC → PTX → SASS               hipcc(LLVM) → AMDGCN ISA

HIP은 CUDA와 의도적으로 거의 같은 모양입니다. 함수 이름의 cuda 접두사가 hip으로 바뀌고, 커널 정의와 실행 문법은 사실상 동일합니다. 이것은 우연이 아니라 설계 목표입니다. 그리고 HIP 코드는 NVIDIA GPU에서도 컴파일됩니다. HIP을 CUDA 위의 얇은 층으로 쓸 수 있다는 뜻이고, 두 벤더를 모두 지원해야 하는 라이브러리가 HIP을 단일 소스로 택하는 근거가 됩니다.

여기서 ROCm 7.14.0에 대해 짚어 둘 것이 있습니다. 버전 번호가 7.2.4에서 7.14.0으로 건너뛴 것이 오타처럼 보이지만 실제입니다. AMD의 릴리스 노트가 "7.9.0 프리뷰에서 시작된 버전 번호 불연속"을 명시하고, 동시에 ROCm이 TheRock이라는 모듈형 빌드 및 릴리스 시스템으로 전환한다고 설명합니다. 핵심 SDK를 가볍게 하고 AI, 데이터 사이언스, HPC용 도메인 SDK를 선택 설치하는 구조입니다.

이 사실 자체가 실무에 주는 정보가 있습니다. ROCm은 지금도 구조가 바뀌고 있습니다. 설치 절차와 패키지 이름이 메이저 릴리스마다 달라질 수 있다는 뜻이고, 문서를 볼 때 버전을 반드시 맞춰 봐야 한다는 뜻입니다. CUDA 쪽은 이 축에서 훨씬 안정적입니다.

라이브러리 대응표

여기서 AMD 특유의 이중 명명을 이해해야 합니다. AMD 공식 문서의 설명은 이렇습니다. roc 접두사 라이브러리는 AMD GPU를 겨냥해 HIP으로 작성된 네이티브 고성능 구현이고, hip 접두사 라이브러리는 CUDA 대응 API를 구현한 이식용 래퍼입니다. hipBLAS는 스스로를 "마셜링 라이브러리"라고 부르며, 뒤에 rocBLAS를 둘 수도 cuBLAS를 둘 수도 있습니다.

고르는 기준은 명확합니다. CUDA에서 옮겨 오는 중이면 hip 쪽을 씁니다. 호출 형태가 같아서 코드 변경이 최소입니다. AMD를 주 타깃으로 새로 쓴다면 roc 쪽을 씁니다. 한 겹이 없고 AMD 전용 기능에 직접 접근합니다.

표의 아래 세 줄은 대응이 한 칸씩 비어 있습니다. cuDNN, NCCL, CUTLASS에는 이식용 래퍼가 없고 별도 이름의 대응물만 있습니다. 이는 API 형태가 달라서 소스 수정 없이는 못 바꾼다는 뜻이고, 딥러닝 프레임워크가 이 층을 각자 흡수해 주기 때문에 대부분의 사용자에게는 문제가 되지 않습니다. 반대로 이 라이브러리들을 직접 호출하는 코드를 들고 있다면, 그 부분이 이식 작업의 핵심 비용이 됩니다.

LLM 추론에 한정하면 최근 등장한 AITER(AI Tensor Engine for ROCm)가 중요합니다. AMD가 LLM 추론용 커널을 모아 둔 저장소이고, vLLM의 ROCm 설치 문서가 빌드 절차에 포함해 두고 있습니다. NVIDIA 쪽의 FlashInfer나 FlashAttention이 차지하는 자리에 해당합니다.

이식 경로 — HIPIFY가 되는 것과 안 되는 것

도구는 두 개입니다.

실무에서는 hipify-perl로 시작하는 경우가 많습니다. 대규모 코드베이스에서 CUDA 헤더를 전부 맞춰 주기가 번거롭기 때문입니다.

# 단일 파일 미리보기 (원본을 건드리지 않고 변환 결과만 출력)
hipify-perl kernel.cu

# 디렉터리 전체를 제자리 변환하고 백업을 남긴다
find . -name "*.cu" -o -name "*.cuh" | xargs hipify-perl -inplace -print-stats

# 정확도가 필요하면 clang 기반
hipify-clang kernel.cu -- -I/usr/local/cuda/include

자동으로 옮겨지는 것

• 런타임 API 호출 이름 (cudaMalloc, cudaMemcpy, cudaStreamCreate 등)
• 커널 한정자와 실행 문법 (__global__, __device__, 삼중 꺾쇠 실행)
• 내장 변수 (threadIdx, blockIdx, blockDim)
• 대부분의 수학 내장 함수
• 헤더 포함문

여기까지가 코드량 기준 90퍼센트쯤 됩니다.

손으로 고쳐야 하는 것

• 인라인 PTX 어셈블리. PTX는 NVIDIA 가상 ISA입니다. AMD에는 대응물이 없습니다. 해당 부분은 HIP 내장 함수로 다시 쓰거나 AMDGCN 인라인 어셈블리로 새로 써야 합니다. 이식 작업에서 가장 확실하게 시간을 먹는 항목입니다.
• 워프 크기 전제. 앞 절의 전부. 자동 변환은 32라는 숫자가 워프 크기인지 다른 뜻인지 알 수 없으므로 손대지 않습니다.
• 워프 수준 프리미티브의 의미 차이. __shfl_sync 계열의 마스크 인자, __ballot의 반환 폭, 명시적 동기화 의미론이 다릅니다. 이름은 옮겨지지만 의미는 검증이 필요합니다.
• CUDA 전용 라이브러리 호출. cuDNN, CUTLASS, cuBLASLt처럼 대응 래퍼가 없는 것들.
• 드라이버 API 사용 코드. cuModuleLoad처럼 저수준 드라이버 API를 직접 쓰는 부분.
• 성능 튜닝 상수 전부. 타일 크기, 블록 크기, 언롤 정도, 파이프라인 단계 수. 변환은 되지만 최적값이 다릅니다. 이식과 최적화는 별개의 작업이고, 여기를 건너뛰면 "돌아가긴 하는데 절반 속도"로 끝납니다.

이식 후 반드시 해야 할 검증

# 1. 수치 검증부터. 성능은 그다음이다.
#    특히 리덕션이 들어간 커널을 집중적으로 본다.
pytest tests/ -k "reduction or attention or norm"

# 2. AMD 쪽 프로파일러로 실제 병목을 다시 잰다.
#    NVIDIA에서의 튜닝 결과를 그대로 믿으면 안 된다.
rocprofv3 --stats -- ./my_app
rocprofv3 --kernel-trace -- ./my_app

Nsight Compute에 해당하는 AMD 도구는 rocprofv3와 ROCm Compute Profiler입니다. 개념은 같습니다. 커널별 시간, 메모리 처리량, 캐시 적중률, 점유율을 봅니다.

현실적으로 무엇이 걸림돌인가

하드웨어가 비슷한데도 현장에서 차이가 나는 이유를 세 가지로 정리합니다.

1. 커널 생태계의 축적량

이것이 가장 큰 항목입니다. 새 모델 구조나 새 양자화 포맷이 나오면, 그것을 위한 커널이 거의 언제나 CUDA로 먼저 나옵니다. FlashAttention의 새 변형, 새 MoE 라우팅 커널, 새 저정밀 GEMM이 모두 그렇습니다. AMD 지원은 몇 주에서 몇 달 뒤에 따라옵니다.

Triton이 이 격차를 실질적으로 줄이고 있습니다. Triton 저장소의 third_party에는 nvidia와 amd 백엔드가 나란히 들어 있어서, Triton으로 쓴 커널은 양쪽에서 컴파일됩니다. 앞선 글에서 다룬 대로 최근의 커스텀 어텐션 커널이 Triton으로 먼저 나오는 경향이 있으므로, 이 경향이 이어지면 격차는 계속 줄어듭니다. 다만 손으로 쓴 CUDA 커널에 의존하는 부분은 여전히 시차가 있습니다.

2. 프레임워크 지원의 시차와 지원 매트릭스

구체적인 예를 하나 보겠습니다. vLLM v0.26.0의 ROCm 설치 문서는 이렇게 적고 있습니다.

• ROCm 6.3 이상을 지원하며, 미리 빌드된 휠은 ROCm 7.0과 ROCm 7.2.1용으로 제공됩니다.
• 지원 GPU는 MI200 계열(gfx90a), MI300(gfx942), MI350(gfx950), Radeon RX 7900 계열, RX 9000 계열, Ryzen AI 계열입니다.
• MI350은 ROCm 7.0 이상을 요구합니다.

여기서 시차가 보입니다. ROCm 최신은 7.14.0인데 vLLM의 사전 빌드 휠은 7.2.1까지입니다. 최신 ROCm을 쓰려면 소스 빌드로 넘어가야 하고, 그러면 Triton, FlashAttention, AITER의 검증된 브랜치를 각각 맞춰 빌드하는 절차가 따라붙습니다. vLLM 문서가 그 브랜치 값들을 Dockerfile에서 확인하라고 안내하는 것 자체가 이 조합의 취약함을 보여 줍니다.

NVIDIA 쪽에서 같은 작업은 대체로 pip install vllm 한 줄입니다. 이 차이가 순수한 소프트웨어 공학적 마찰이고, 팀의 시간을 먹습니다.

실무적 결론은 명확합니다. AMD에서는 검증된 컨테이너 이미지를 쓰는 것이 사실상 기본 경로입니다. vLLM은 공식 이미지 vllm/vllm-openai-rocm을 Docker Hub에 올려 두고 있으며, AMD가 배포하던 rocm/vllm 계열 이미지는 이제 공식 이미지 쪽으로 넘어갔습니다. 직접 빌드하겠다는 결정은 그 자체로 상당한 유지 비용을 부르는 결정입니다.

3. 문제가 생겼을 때의 정보량

측정하기 어렵지만 실재하는 항목입니다. 오류 메시지를 검색했을 때 나오는 결과의 양, 스택 오버플로 답변의 수, 같은 문제를 겪은 사람의 블로그 글, 이미 답이 나와 있는 GitHub 이슈. 이 모든 것이 CUDA 쪽이 압도적으로 많습니다.

디버깅 시간은 프로젝트 일정에 그대로 들어가는 비용입니다. 하드웨어 가격표에는 안 나오지만 총소유비용에는 들어갑니다.

어떤 워크로드라면 AMD가 합리적인가

위 걸림돌들을 인정한 뒤에도 AMD가 합리적인 경우가 분명히 있습니다. 조건은 구체적입니다.

첫째, 널리 쓰이는 모델을 추론으로 서빙하는 경우. Llama, Qwen, DeepSeek 계열의 주류 모델을 vLLM으로 서빙하는 일은 이미 잘 다져진 경로입니다. 검증된 컨테이너를 쓰고 표준 모델을 올리는 한, 위에서 말한 마찰의 대부분은 남이 이미 겪고 해결해 둔 것입니다.

둘째, 메모리 용량이 결정적인 경우. 큰 모델을 더 적은 수의 GPU에 올릴 수 있으면 텐서 병렬 통신이 줄고, 그 자체로 성능과 단순함을 얻습니다. AMD가 세대별로 용량에서 앞서는 구간이 있어 왔고, 그 구간에서는 이것이 실질적인 이점입니다.

셋째, 조달과 가격 협상력이 필요한 경우. 대량 도입에서 두 번째 공급자의 존재는 그 자체로 가치입니다. 실제로 쓰지 않더라도 협상 테이블에 올릴 수 있는 대안이 있는 것과 없는 것은 다릅니다.

넷째, 스택 상단에서만 일하는 팀. PyTorch 위에서만 작업하고 커스텀 CUDA 커널을 들고 있지 않다면, 이식 비용의 큰 부분이 애초에 없습니다.

반대로 AMD를 피해야 하는 경우도 명확합니다.

• 손으로 쓴 CUDA 커널이나 인라인 PTX에 성능을 의존하는 코드베이스가 있는 경우
• 새 모델 구조와 새 커널을 남보다 먼저 써야 하는 연구 조직
• GPU 인프라에 전담 인력을 둘 수 없는 소규모 팀
• cuDNN이나 CUTLASS를 직접 호출하는 코드가 핵심에 있는 경우

한 가지 실무 조언을 덧붙이면, 지금 CUDA로 새 커널을 쓰고 있다면 Triton으로 쓸 수 있는지부터 검토하는 것이 가장 값싼 보험입니다. 나중에 AMD를 검토하게 되었을 때, Triton 커널은 재컴파일로 끝나고 CUDA 커널은 이식 프로젝트가 됩니다.

마치며 — 격차는 실리콘이 아니라 축적된 커널에 있다

두 회사의 GPU는 같은 세대에서 비슷한 연산 능력과 비슷한 메모리 대역폭에 도달해 있습니다. SM과 CU는 개념적으로 같고, 텐서 코어와 매트릭스 코어도 하는 일이 같습니다. HIP은 CUDA와 거의 같은 API를 제공하고, HIPIFY는 코드의 90퍼센트를 자동으로 옮겨 줍니다. 여기까지만 보면 차이가 없어 보입니다.

차이는 그 아래에 쌓인 것에 있습니다. 지난 15년간 CUDA를 겨냥해 작성되고 튜닝된 커널의 총량, 그 커널들을 전제로 만들어진 라이브러리, 그 라이브러리를 전제로 쓰인 프레임워크, 그리고 그 전부에 대해 축적된 디버깅 지식. 이것이 하루아침에 복제되지 않는 자산이고, 실제 성능 격차의 대부분을 설명합니다.

그리고 이 격차가 줄어드는 방향도 여기서 나옵니다. 커널이 특정 벤더의 언어가 아니라 Triton 같은 이식 가능한 층에서 작성될수록, 축적의 효과는 한쪽에만 쌓이지 않습니다. 컴파일러가 벤더 종속성을 흡수하는 층이 두꺼워질수록 선택은 실제로 자유로워집니다. 그것이 이 시리즈의 다음 글에서 다룰 주제입니다.

지금 판단해야 한다면 질문은 하나로 줄어듭니다. 우리 성능이 우리가 직접 쓴 CUDA 코드에 걸려 있는가, 아니면 남이 쓴 라이브러리에 걸려 있는가. 후자라면 AMD는 검토할 만한 선택지입니다. 전자라면 이식 비용을 먼저 정직하게 견적해야 합니다.

참고 자료

• ROCm 공식 문서: https://rocm.docs.amd.com/
• ROCm 릴리스(버전과 날짜 확인): https://github.com/ROCm/ROCm/releases
• HIP 포팅 가이드(워프 크기 경고 포함): https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html
• ROCm API 라이브러리 목록(CUDA 대응 관계): https://rocm.docs.amd.com/en/latest/reference/api-libraries.html
• HIPIFY 저장소와 문서: https://github.com/ROCm/HIPIFY
• AITER (AI Tensor Engine for ROCm): https://github.com/ROCm/aiter
• vLLM ROCm 설치 문서: https://docs.vllm.ai/en/latest/getting_started/installation/gpu.html
• AMD MI300X 튜닝 가이드: https://rocm.docs.amd.com/en/latest/how-to/tuning-guides/mi300x/index.html
• AMD Instinct 제품 페이지(공식 스펙): https://www.amd.com/en/products/accelerators/instinct.html
• NVIDIA 데이터센터 GPU 제품 페이지: https://www.nvidia.com/en-us/data-center/