지난 달에 나온 deepseek는 최고 티어 AI와 비견되는 성능에 1/10 이하의 학습 비용 및 저렴한 장비 사용으로 많은 관심을 받고 있어.
deepseek를 그냥 알고만 있다가 기술적으로 어떻게 최적화했는지 소개하는 좋은 글을 보게 되어서 일부 내용과 함께 그 부분에 대한 지식을 정리해봤어.
아래 원본 글에는 최적화 관련 내용 외에 실제 AI 학습 전략이나 다른 정보들도 많으니 한번 읽어보면 재밌을거야.
https://stratechery.com/2025/deepseek-faq/
https://dev.to/datamonk_/how-deepseek-is-making-high-performance-ai-accessible-to-all-26fp


- 배경
deepseek가 저렴한 장비를 사용할 수 밖에 없던 이유는 다들 알다시피 미중갈등으로 인해 미국이 중국을 대상으로 수출 규제를 시작했기 때문이야. 그 대상에는 반도체 설비 기기 외에 고성능 gpu도 포함되었어. 그래서 AI학습에 필수적인 nvidia의 a100, h100 등이 2022년 8월부터 수출이 제한되었고, 이는 중국에게도 큰일이지만 nvidia 회사로 봤을 때에도 큰 매출 손실이 나는 부분이었지.
그래서 nvidia에서는 h100의 저가형 모델(=중국 전용)인 h800을 만들어서 수출했는데 이조차 2023년 10월에 막히게 돼. 지금은 그보다 낮은 성능의 h20 모델을 만들고 수출 중인데 이것도 언제 막힐지 모르지 뭐. 성능은 대충 절반씩 깎인다고 보면 될거야. (어떤 기사에 따르면 h100에 비해 h20은 1/5 수준의 성능이라고..)
그 사이에 중국이 얼마나 수입했는지 궁금해서 기사 몇개 찾아보니까 2022년 중국 매출이 50억달러였는데 2023년에는 100억달러가 찍혔더라. 2024년은 아직 발표하지 않았지만 최소한 전년보단 많다고 하네.

중국 입장에서 AI 개발을 하려니, 저성능 gpu도 문제지만 반토막난 대역폭도 큰 문제였어. 왜냐하면 AI 학습엔 파라미터가 수천억개라서 하나의 gpu 메모리로는 부족하여 여러 개의 gpu를 묶어서 사용할 필요가 있거든. 안그래도 느린 성능인데 대역폭도 절반이니 이대로는 AI 개발이 쉽지 않았던거야.
이 한계를 극복하기 위해 AI 전략의 변경과 함께 low-level 최적화를 시도했는데, PTX로 여러 방면의 최적화를 적용하여 deepseek를 개발하게 돼.


- PTX란?
PTX는 Parallel Thread Excution의 약자로, nvidia에서 만든 ISA (Instructure Set Architecture)야. cpu에서는 어떤 언어든 컴파일해서 어셈블리어로 변경 후 사용하는 것처럼, gpu도 cuda 코드(c, python..)를 그대로 사용하는게 아니라 그걸 컴파일하면 ptx 중간 언어로 변환되고, 이를 gpu가 읽으면 드라이버가 기계어로 변환 후 사용하게 돼. (다만 ptx는 nvidia gpu에 종속된 언어라 amd는 안됨)
c/c++을 사용하다가 인라인 어셈을 사용하듯이 cuda코드에도 사용 중간에 ptx코드를 넣을 수 있다는 거지.


- PTX 최적화 방식
low-level 최적화를 시도할 수 밖에 없던 배경과 ptx가 뭔지 알아봤으니 이제 deepseek팀에서는 이걸로 뭘했는지 자세히 알아볼게.

1. cuda코드를 컴파일하면서 비효율적인 계산이 발생하던 부분을 미세 조정을 통해 효율적으로 사용되도록 재구성
2. tensor 연산을 재구성하여 메모리 병목 현상을 줄이고 처리량 상승
3. ptx의 핵심 커널을 다시 작성하여 처리량 상승

여기서 1,2번은 그럴 수 있다 치고 3번을 조금 더 깊이있게 볼게.


- gpu 구성 및 최적화 방식
일단 gpu가 어떻게 구성되어 있는지 알아야하니 아래 h800의 스펙을 한줄씩 해석해보면,

* Shading Units: 16896
  병렬 연산을 처리하는 코어 개수
* TMUs: 528
  Texture Mapping Unit. 모델에 이미지를 입히는 가속화 코어
* ROPs: 24
  Raster Operations Processor. 렌더링된 최종 이미지의 깊이 테스트, 프레임 버퍼 쓰기 등의 작업 수행
* SM Count: 132
  Streaming Multiprocessor. Shading Unit과 TMU, L1 Cache 등을 포함하는 독립적인 연산 단위
* Tensor Cores: 528
  딥러닝에 특화된 행렬 연산의 가속화 코어
* L1 Cache: 256 KB (per SM)
  SM마다 하나씩 할당된 캐시 메모리
* L2 Cache: 50 MB
  모든 SM에서 공유하는 2차 캐시 메모리

우리가 보통 cuda 코어개수라고 부르는게 Shading Unit(이하 SU)인데 저 많은 16896개가 개별적으로 작동하진 않고, 하나의 SM(Streaming Multiprocessor)마다 SU(128개), TMU(4개), Tensor(4개)가 포함된 구성 요소인거야.
여기서 h800은 132개의 SM 단위로 이루어져 있고, SM마다 여러 개의 warp를 실행할 수 있는데 warp 하나마다 32개의 스레드를 실행할 수 있어. 그리고 h800에서 SM 하나가 실행 가능한 warp 개수는 64개이고, SM 하나당 최대 2048개의 스레드를 동작시킬 수 있으니 h800 하나가 이론상 27만개가 넘는 스레드가 동시에 실행이 가능해.

이게 대충 머릿속에 그려졌다면 이제 이해하기 쉬울거야.
h800에서는 대역폭이 낮다고 했잖아? 그래서 deepseek팀은 저 132개의 SM중에 20개를 외부 칩과의 통신을 관리하도록 별도로 구성하여 연산 결과가 지연없이 데이터 전송되는 최적화로 낮은 대역폭을 극복했어. 이는 cuda에서는 불가능한 기능이고 ptx를 통해서만 가능한 최적화였던거지.

그리고 h800은 연산 성능과 대역폭이 낮기 때문에 데이터를 8비트짜리 fp8포맷인 e4m3(지수부 4비트, 가수부 3비트)를 사용하여 연산 속도 상승과 데이터 양 감소 효과를 얻을 수 있었어. h800은 fp64계산이 끔찍하게 느린데 반해 fp8/16은 h100과 비슷했거든. 참고로 e4m3의 정밀도는 0.125이고 최대 32768에서 최소 0.000061까지 표현할 수 있어.


- 논문 살펴보기
아티클이나 기사만으로는 한계가 있어서 공개된 논문에서 ptx 부분만 살짝 봤는데 생각보다 꽤 자세한 내용이 있었어.
https://github.com/deepseek-ai/DeepSeek-V3/blob/main/DeepSeek_V3.pdf
주요 챕터는 Efficient Implementation of Cross-Node All-to-All Communication 이었는데 여기에서 필요한 내용만 요약해볼게.

1. (h800의 NVLink 개수는 최대 8개이고 이만큼의 그래픽 카드를 NVLink로 연결한 걸 하나의 노드라고 정의함) 노드 내 통신은 NVLink로 연결. 대역폭은 160GB/s
2. 클러스터(노드-노드간) 연결은 IB(InfiniBand)로 연결되어 있고, 대역폭은 50GB/s
3. IB가 NVLink보다 상대적으로 느리기 때문에 질의가 노드를 벗어나지 않도록 토큰을 최대 4개 노드에 중복하여 저장하여 최대한 노드 내에서 검색되도록 하고, 각 토큰은 대상 전문가에게 즉시 전달 (여기는 deepseek AI 전략인 MoE를 알아야 하니 여기서는 생략)
4. 20개의 고정된 SM을 10개의 통신 채널로 나눠서 듀얼 파이프라이닝 구성
5. warp 특수화 처리로 디스패칭 및 결합 과정을 처리하고, 각 통신 작업에 할당된 warp 수를 동적으로 조정
6. 디스패칭 및 결합 과정이 계산 스트림과 겹치도록 설계하여 유휴시간이 최소화되도록 구성 (cpu 파이프라이닝 RAW hazard 해결책 느낌?)
7. 데이터에 맞춤형 PTX 명령셋을 제공하고 통신 청크 크기를 자동 조정하여 L2 캐시 사용량이나 다른 SM에 대한 간섭 감소

warp 개념은 설명을 안했는데 어차피 구글링하면 그림과 함께 설명한 내용이 많으니 여기서는 생략할게.


- PTX 살펴보기
실제로 어떤 ptx 명령으로 최적화했는지 궁금해서 ptx 명령셋을 살짝 구경해봤어. 
ptx는 nvidia에서 만든 ISA라서 메뉴얼이 상당히 잘 되어있는 편이야. 아래 링크에서 구경해 봐.
https://docs.nvidia.com/cuda/parallel-thread-execution/

일단 기본적인 느낌은 어셈블리어와 상당히 유사한데 명령셋이 보기 편하게 구성되어 있더라. 예를 들어 mov를 보면 어셈블리어에서는 복사할 타입마다 명령어가 따로 있는데 (movl, movq..) ptx에서는 mov.u8, mov.f32로 명령 뒤에 타입을 붙여서 구성하고 있어. add, sub같은 다른 명령도 동일해 (add.s32, sub.f16..)

deepseek에서 적용한 최적화 기법인 별도로 SM id를 지정하여 실행하는 명령이 있는지 쭉 살펴봤는데 그런건 없더라고. 생각해보니 이 코드를 실행하는 상황이라면 이미 SM 안쪽일거라서 다른 SM으로 옮겨서 실행한다는게 말이 안되긴 해.
그래서 다시 찬찬히 보니까 %smid 라는 특수 레지스터에서 현재 실행하는 SM id를 얻어올 수 있었고, 이를 통해 132개의 SM중에 상위 20개만 제어하는 코드를 만들 수는 있겠더라.

이걸 대충 ptx 코드로 써보면,

// 32비트 변수 id를 정의
.reg    .s32 id;
// SM id를 id 변수로 복사
mov.u32 id, %smid;
// id가 20보다 작으면 p에 true 저장 (0~19)
setp.lt.u32 p,id,20
// p가 true면 L1 레이블로 이동
@p  bra L1
// L1 레이블에서 개별 SM id 데이터에 따라 연산된 데이터의 최적화 진행
L1:
...

실제로 이렇게 했는지는 모르겠지만 명령셋 내에서는 이거 말고는 아이디어가 안떠오르네. deepseek가 코드 공개했다고 해서 그 repo도 가봤는데 ptx 관련 로직은 없는걸 보니 핵심 기술이라 비공개인듯.
smid 외에 warpid, tid, gridid도 얻어올 수 있고, async 명령도 존재해서 스레드 제어도 가능하니 이를 조합해서 사용하지 않았을까 싶어. (async 코드는 쓰다보니 지나치게 길어져서 생략함)


- 잡설 및 소회
국내에는 펌글만 많고 정보가 없어서 해외 기준으로 많이 찾아봤는데 글을 마무리하다가 네이버에서 양질의 좋은 기사를 하나 봤어.
https://naver.me/5mIoUqAX
여기서 말하길, deepseek팀은 이전에 금융권의 HFT(High-frequency trading, 고빈도 매매)를 개발하던 팀이라 대역폭이나 지연 시간 최적화에 노하우가 많은 팀이었다고 해.
그걸 알고나서 보니 이게 혜성처럼 나타난 기술은 아니고 숙련된 팀의 기술력으로 이뤄낸 성과라는게 참 생각할 부분이 많았어. 이런 세계적인 성과들은 개인의 기술력이 아닌 팀 단위의 기술력이 꼭 필요하구나 싶고..
nvidia의 cuda외에 amd의 ROCm이랑 HIP도 같이 다뤄보려고 며칠동안 살펴봤는데 아직 글로 쓸 정도의 깊이까진 도달하지 못해서 나중에나 써봐야지 싶네.


- 결론
1) deepseek는 저성능 하드웨어를 사용할 수 밖에 없는 환경의 극복을 위해 많은 최적화가 이루어졌다
2) nvidia 그래픽 드라이버의 중간 언어인 PTX를 사용하여 low-level 최적화로 낮은 대역폭을 극복했다
3) PTX는 어셈블리어와 비슷하고 async 명령어로 스레드 제어가 가능하다
4) 중국의 low-level 기술력은 참 대단한 느낌


이전글: https://frogred8.github.io/
#frogred8 #deepseek #ptx

---