― 알렉스넷 원본 소스코드를 직접 파 내려간, 어느 아마추어 AI 사학자의 발굴 노트
이수화 · 서울대 빅데이터혁신융합대학 연구교수 · 법무법인 디엘지 AI센터장
“일리야는 하자고 했고, 알렉스가 그걸 돌아가게 만들었고, 상은 내가 받았다.”
— 제프리 힌튼, 2024년 노벨상 수상자
앞선 글에서 나는 2012년의 기적을 ’야생의 요람’이라 불렀다. 좋은 반응에 우쭐하던 차, 문득 못된 버릇이 도졌다. 역사학자는 결국 사료(史料) 앞에서만 정직해진다. 남의 요약을 인용하는 자는 논평가일 뿐, 원본을 만지는 자라야 (아마추아) 사학자다. 그래서 나는 삽을 들고 지층을 파 내려가기로 했다. 마침 미국 컴퓨터역사박물관(CHM)이 구글과 5년을 협상해 2025년 3월 세상에 풀어놓은 그 코드 ― github.com/computerhistory/AlexNet-Source-Code ― 가 나를 기다리고 있었다.
미리 고백하자면, 이 발굴에는 함정이 하나 있었다. 얼마 전 나는 어느 생성형 AI에게 “알렉스넷 원본 코드를 보여 달라”고 청했다가, 그럴듯한 위작(僞作)을 한 아름 받아 들었다. filter_acts.cu라는 진짜 파일명에, 진짜처럼 생긴 커널 코드를 얹은 정교한 가짜였다. 기계는 사료를 ‘지어낼’ 줄 안다. 그러니 더더욱, 역사학자는 제 손으로 저장소를 열어야 한다. 아래에 인용하는 코드는 단 한 줄도 상상이 아니다. 전부 그 저장소에 실재하는, 2012년 여름의 화석들이다.
1. 물증(物證): 3,072메가바이트짜리 감옥
저장소를 열면 코드보다 먼저 눈에 걸리는 파일이 하나 있다. deviceQuery.txt. 코드가 아니라, 기계가 제 몸을 스스로 진술한 조서(調書)다. 첫 줄의 경로부터 냄새가 난다.
/u/kriz/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/deviceQuery Starting... Found 4 CUDA Capable device(s) Device 0: "Tesla S2050" CUDA Capability Major/Minor version number: 2.0 Total amount of global memory: 3072 MBytes (3220897792 bytes) (14) Multiprocessors x ( 32) CUDA Cores/MP: 448 CUDA Cores Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 32768 Warp size: 32
/u/kriz/. 알렉스 크리제프스키의 홈 디렉터리다. 십수 년 전 토론토의 한 사용자 계정이, 지문처럼 파일 첫머리에 찍혀 있다. 그런데 조서를 조금만 더 읽으면 앞뒤가 맞지 않는 대목이 나온다. Found 4 CUDA Capable device(s), 그리고 Device 0: "Tesla S2050". 논문이 밝힌 학습 장비는 분명 GeForce GTX 580 두 장이었는데, 이 진술서는 GTX가 아니라 테슬라(Tesla)를, 그것도 두 장이 아니라 네 장을 가리킨다. 여기서 대충 얼버무리면 독자만 어리둥절해진다. 물증은 끝까지 캐물어야 물증이다.
열쇠는 448이라는 숫자에 있다. GTX 580은 512개의 CUDA 코어를 완전히 개방한 게이밍 최고급 카드다. 반면 조서에 찍힌 건 448 CUDA Cores — 같은 페르미(GF100) 칩에서 일부 연산 유닛을 잠근 채 안정성을 택한 기업용 등급, 곧 테슬라 M2050의 지문이다. 그리고 Tesla S2050은 데스크톱 카드 한 장이 아니라, 이 M2050 네 개를 1U 랙 섀시 하나에 몰아넣은 서버다. deviceQuery가 장치를 4개로 세는 건 그래서 당연하다. 여기까지는 스펙이 말해주는 물증이다.
그렇다면 이 로그는 크리제프스키의 방이 아니라, 대학 연구실 어딘가의 랙 서버 위에서 찍혔다고 보는 편이 자연스럽다. /u/kriz/라는 공용 유닉스 계정의 냄새까지 겹치면, 짐작건대 힌튼 연구실이 함께 쓰던 서버였을 것이다. 그렇다면 이 조서는 AlexNet의 최종 산실이 아니라, 그 이전 개발·테스트 단계에 남은 화석일 공산이 크다. 값비싼 공용 서버는 대학원생 하나가 일주일 내내 독점할 수 있는 자원이 아니었을 테니, 뼈대는 이 공용 서버에서 세우고 닷새 넘게 걸리는 최종 학습은 500달러짜리 게이밍 카드 두 장을 따로 구해 방구석 PC에 꽂아 끝까지 혹사시켰으리라. 여기서부터는 물증이 아니라 정황이다. 다만 흩어진 단서를 잇는 가장 그럴듯한 동선이기도 하다.
그런데 진짜 주인공은 장비의 종류가 아니다. 조서의 한 줄, Total amount of global memory: 3072 MBytes. 3기가바이트. 공용 테슬라의 칩 하나도, 크리제프스키의 GTX 580 한 장도, 똑같이 이 3GB 페르미의 벽 앞에 갇혀 있었다. 감옥은 기계를 가리지 않았다. 세상을 바꾼 신경망은 딱 이 3,072메가바이트라는 감옥 안에서 태어났고 — 더 정확히는, 그 감옥이 신경망을 낳았다. 한 장에 다 담기지 않는 모델을 정확히 반으로 쪼개 두 GPU에 나눠 싣고 특정 층에서만 서로 말을 섞게 한 병렬의 묘수는, 메모리가 넉넉했다면 태어나지 않았을 발명이었으니까.
2. 연장(道具): 맨손으로 벼려낸 커널
역사는 종종 무기가 아니라 연장을 벼린 자를 잊는다. AlexNet을 이야기할 때 우리는 으레 ‘신경망 구조’를 떠올리지만, 정작 2012년의 진짜 사건은 아이디어가 아니라 연장에서 벌어졌다. 합성곱도 역전파도 이미 수십 년 전에 나온 개념이었다. 없던 것은 그 수학을 실제로, 빠르게, 거대한 규모로 굴려낼 ‘기계를 부리는 법’이었다. 이 글 맨 앞에 걸어둔 힌튼의 회고에서, 정작 우리가 오래 들여다봐야 할 대목은 상을 받은 세 번째 토막이 아니라 그 가운데다. 알렉스가 그걸 돌아가게 만들었다 — 딥러닝 혁명의 뇌관은 바로 그 ‘돌아가게 만들었다’는 한마디 안에 있다.
돌아가게 만든다는 건 무슨 뜻인가. 여기서 CUDA와 병렬처리를 짚고 가야 한다. CPU는 소수의 박사들이다. 똑똑하지만 몇 명뿐이라 문제를 한 줄씩 순서대로 푼다. GPU는 정반대다. 산수만 겨우 하는 수천 명의 아이를 한 방에 몰아넣고 똑같은 곱셈을 동시에 시키는 물량 작전이다. 신경망 학습이란 결국 똑같은 곱셈-덧셈을 수백만 번 되풀이하는 일이라, 이 물량 작전과 궁합이 완벽하다. 엔비디아의 CUDA(2007)는 원래 그래픽을 그리는 데만 쓰이던 이 부대를 일반적인 수학 계산에 동원할 수 있게 열어준 지휘 언어였다.
문제는 궁합이 좋다는 것과 실제로 빠르게 돈다는 것이 전혀 다른 이야기라는 데 있다. 수천 명을 아무렇게나 부리면 절반은 놀고 절반은 서로 부딪힌다. 진짜 속도는 누가 언제 메모리의 어느 칸을 건드리는지까지 손으로 짜 맞출 때 비로소 나온다. 어느 데이터를 레지스터에 얹고, 어느 것을 공유 메모리에 잠깐 세워두고, 3GB뿐인 VRAM 예산을 어떻게 쪼갤지 — 이건 통계학자의 일이 아니라 하드웨어 대장장이의 일이다. 그리고 2012년에는 이 일을 대신 해줄 연장이 없었다. 파이토치도 텐서플로도 없었고, 엔비디아가 딥러닝용 최적화 부품을 담아 내놓은 cuDNN조차 2년 뒤인 2014년에야 나왔다. 실험을 돌리기 전에 실험 도구부터 맨손으로 만들어야 했다는 뜻이다. 대부분의 연구실이 이 통행세 앞에서 멈춰 선 이유다. 신경망을 아는 머리와 GPU를 바닥까지 훑는 손을 한 사람 안에 가진 경우가, 그 시절엔 거의 없었다.
크리제프스키는 그 드문 한 사람이었다. 공개된 소스를 열면 그가 치른 대가가 파일 이름에 그대로 박혀 있다. 핵심은 파이썬 스크립트 몇 줄이 아니라 cuda-convnet — 그가 밑바닥부터 손으로 쓴 C++/CUDA 라이브러리다. conv.cu 같은 파일 안에는 합성곱을 GPU에서 돌리는 커널이 직접 벼려져 있고, 층의 구조와 연결은 options.cfg 설정 파일 한 장으로 갈아 끼울 수 있게 설계돼 있다. 그리고 백미는, 3GB 한 장에 다 담기지 않는 모델을 정확히 반으로 갈라 두 GPU에 나눠 싣고 특정 층에서만 데이터를 교환하게 만든 ‘모델 병렬(model-parallelism)’ 코드다. 1번에서 본 감옥의 벽을, 그는 벽을 부수는 대신 방을 둘로 쪼개는 방식으로 통과했다. 포크된 저장소 한 귀퉁이에는 “비어 있는 GPU 보드를 찾아 잠그는 스크립트”를 호출하는 대목까지 남아 있는데, 이것이야말로 여럿이 한 대의 랙 서버를 눈치껏 나눠 쓰던 대학 연구실의 체취다 — 1번의 Tesla S2050 로그와 정확히 같은 세계다.
이 연장이 얼마나 잘 벼려졌는지는, 훗날 남들이 그것을 어떻게 대했는지가 증언한다. 이후 여러 연구자와 프레임워크가 자기 신경망을 새로 짜는 대신 크리제프스키의 커널을 통째로 가져다 썼고, 어떤 코드에는 “알렉스 덕분에 엄청난 시간을 아꼈다. 무한히 감사한다”는 주석까지 달렸다. 세상을 바꾼 건 어떤 우아한 방정식이 아니라, 그 방정식이 실제로 굴러가게 만든 이 땀내 나는 연장이었던 셈이다. 컴퓨터역사박물관이 5년을 협상해 끝내 GitHub에 올린 것도 매끈한 이론이 아니라 바로 이 손때 묻은 코드였다. 딥러닝 혁명의 첫 삽은, 수식이 아니라 컴파일러 앞에서 밤을 새운 한 대학원생의 손끝에서 떠졌다.
코드를 열다: 침실에서 나온 네 개의 사료
지어낸 코드를 인용할 필요는 없다. CHM이 공개한 저장소를 직접 열면 된다. 파일 첫머리엔 서명이 박혀 있다. Copyright (c) 2011, Alex Krizhevsky. 2011년 — 논문보다 한 해 앞선 이 날짜가, 연장이 실험보다 먼저 벼려졌음을 말한다. 그 아래에서 네 조각을 꺼낸다.
사료 ① — 3GB 감옥을 통과하는 법: 벽을 부수지 않고, 방을 둘로 쪼갠다
# layers-120.cfg (2012년 우승 설정) [conv1a] type=conv inputs=data filterSize=11 stride=4 filters=32 gpu=0 [conv1b] type=conv inputs=data filterSize=11 stride=4 filters=32 gpu=1
흔히 “한 줄로 GPU를 쪼갰다”는 식으로 낭만화되지만, 실제는 더 우직하다. 대회를 제패한 그 설정 파일(layers-120.cfg)을 열면, 크리제프스키는 신경망 전체를 a와 b 두 기둥으로 통째 복제한 뒤 층마다 gpu=0·gpu=1로 물리적 카드를 못박아 두었다. 왼쪽 절반은 첫 번째 GTX 580에, 오른쪽 절반은 두 번째 카드에. 두 기둥은 정해진 층에서만 서로 말을 섞는다. 3GB라는 벽을 부술 수 없으니, 뇌를 반으로 갈라 방 두 개에 나눠 재운 것이다. 결핍이 설계를 낳은 순간이, 이 한 장의 설정 파일에 박제돼 있다.
사료 ② — 학습의 심장, 그리고 세 사람의 지문
_weightsInc->add(*_weightsGrad, _mom, _lrs->getRate(progress)); // 모멘텀: V ← 0.9·V + ε·∇W _weightsInc->add(*_weights, -_wc * _lrs->getRate(progress)); // 가중치 감쇠: V ← V − ε·λ·W _weights->add(*_weightsInc); // 갱신: W ← W + V
그리고 같은 파일 위쪽, 실험용 갱신 커널과 그 분기에는 세 사람의 지문이 주석으로 남아 있다.
// superMomUpdate 커널 본문 — 모멘텀 0.9가 코드에 그대로 박혀 있다
V[t] = /*mom*/0.9 * V[t] + eps * __fdividef(g[t], sqrtf(G2[t] - G[t] + superEps));
...
superMomUpdate<<<blocks, threads>>>(...);
// Geoff version
} else {
...
// Ilya version
오늘날 optimizer.step() 한 줄에 감춰진 학습의 심장이, 여기서는 날것의 행렬 연산 세 줄로 드러난다. 첫 줄은 모멘텀(관성 0.9로 이전 방향을 기억), 둘째 줄은 가중치 감쇠(과적합을 막는 감가상각), 셋째 줄은 그 결과를 실제 가중치에 더하는 갱신이다. 논문에 적힌 그 유명한 갱신 식이, 자동 미분도 최적화 객체도 없이 손으로 옮겨져 있다. 더 뭉클한 건 같은 파일의 실험용 커널이다. 모멘텀 값 0.9가 코드에 그대로 박혀 있고, 바로 위 블록에서 보듯 지워지지 않은 주석 두 줄이 화석처럼 남아 있다. // Geoff version 그리고 // Ilya version. 갱신 규칙을 두고 세 사람이 각자의 안을 시험하던 흔적이다. 훗날 한 사람은 노벨상을, 한 사람은 OpenAI를 향해 갈라지지만, 이 파일 안에서 그들은 아직 같은 코드를 만지는 세 명의 동료다.
사료 ③ — 오류역전파의 방아쇠, 단 한 줄
const int label = int(labels[tx]); float v = gradCoeff * ((label == ty) - y_l[tidx]); dE_dx_l[tidx] = v;
학습이란 결국 “틀린 만큼 고치는” 일이다. 그 ‘틀린 만큼’이 태어나는 자리가 바로 이 한 줄이다. (label == ty)는 이 출력이 정답이면 1, 아니면 0 — 즉 정답표다. y_l은 신경망이 내놓은 예측 확률이다. 둘을 빼면 정답과 예측의 차이, 곧 오차가 나온다. 소프트맥스와 교차엔트로피라는 거창한 이름의 미분이, 실전에서는 이렇게 뺄셈 한 번으로 압축된다. 이 한 줄에서 태어난 오차 신호가 신경망의 모든 층을 거꾸로 거슬러 오르며 6천만 개의 가중치를 조금씩 밀어낸다. 세상을 바꾼 학습의 방아쇠는, 놀랍도록 단순한 뺄셈이었다.
사료 ④ — 우승본의 지문: 설정 파일이 논문과 악수하다
# layer-params-120-2012.cfg (발췌) [conv3a] epsW=0.00001 # 학습률 momW=0.9 # 모멘텀 wc=0.0005 # 가중치 감쇠 [conv1a] epsW=0.0000 # ← 이 층은 학습률 0, 즉 '얼려둔' 상태
저장소에는 100번대부터 180번대까지, 크리제프스키가 돌린 수십 개의 실험 설정이 지층처럼 쌓여 있다. 그중 2012년 대회의 그 2-GPU 망에 대응하는 것이 120번 실험이다. 근거는 숫자다. 이 파일의 모든 층에 적힌 momW=0.9와 wc=0.0005는, 논문이 명시한 학습 설정(모멘텀 0.9, 가중치 감쇠 0.0005)과 한 치도 다르지 않다. 그리고 그 두 값은 앞서 사료 ②에서 가중치를 밀어 올리던 weights.cu의 바로 그 손잡이다. 논문의 문장과 설정 파일의 한 줄과 GPU 위의 커널이, 같은 두 숫자로 물려 돌아간다. 더 인간적인 대목은 epsW=0.0000이다. 앞쪽 층의 학습률을 0으로 얼려 두고 뒤쪽만 미세조정하던, 학습 한복판의 스냅샷 — 완성된 청사진이 아니라 손을 담근 채 멈춘 작업의 단면이다. 세상을 바꾼 그 망은, 이렇게 얼리고 녹이기를 반복하던 어느 밤의 설정 파일 안에 화석으로 남아 있다.
3. 개고생의 지질학 ― 160개의 실패한 설정 파일
여기서 잠깐, 사학자의 눈이 번뜩이는 대목을 짚고 싶다. layers/ 폴더에는 설정 파일이 하나도, 둘도 아니고 160개가 쌓여 있다. layer-params-100.cfg, -106, -107… -184까지, 번호를 매겨가며. 그중 17개엔 4gpu가, 28개엔 2009·2010·2011·2012 같은 연도가 꼬리표처럼 달렸다. 설정 파일만 모아도 4만 5천 줄이 넘는다.
이 숫자가 바로 내가 하고 싶은 이야기의 핵심이다. 저 160개의 파일은 160번의 시행착오이고, 그 대부분은 실패의 묘비다. layers-141-2009-half.cfg와 layers-141-2009-2012.cfg가 나란히 누워 있는 풍경은, 같은 실험을 조금씩 비틀어가며 수십 번 되돌린 사람의 손때다. 최초로 융합을 시도하는 자에게 우아한 지름길 따위는 없다. 게임용 그래픽 칩과 낡은 신경망 이론이라는, 서로 만날 이유가 없던 두 세계를 처음 포개는 일은 원래 이렇게 생겼다. 오래 걸리고, 지저분하고, 대부분 망한다. 그리고 망한 파일 160개를 버리지 않고 번호까지 매겨 남겨 두는 그 미련함이, 161번째의 기적을 낳는다.
명심할 것. 이 코드가 위대한 이유는 코드가 깨끗해서가 아니다. 정반대다. 이 코드는 더럽기 때문에 위대하다.
4. 젠슨 황의 어려운 선물 ― 손으로 깎은 CUDA 엔진
이 저장소의 언어 분포를 GitHub는 이렇게 집계한다. Cuda 58.1%, Python 39.1%, Shell 1.7%, C++ 1.1%. 앞서 어떤 기계는 “CUDA C++ 80%, 파이썬 15%”라 우겼지만, 진실은 파이썬의 비중이 훨씬 크다(정직한 발굴은 늘 이렇게 소소하게 통념을 배신한다). 그럼에도 무게중심은 분명 CUDA에 있다. 전체 8,449줄의 .cu/.cuh 코드. 젠슨 황이 2006년부터 게이머들의 야유를 들어가며 칩에 욱여넣었던 그 범용 병렬연산 언어를, 크리제프스키는 신경망을 위한 행렬 곱셈기로 개조했다.
문제는, CUDA가 젠슨 황의 기대만큼 순순히 뻗어주는 언어가 아니었다는 데 있다. 어렵고, 까다롭고, 손이 많이 갔다. 얼마나 손이 많이 갔는지는 빌드 스크립트 build.sh 한 장이 다 말해 준다.
#!/bin/sh # 합성곱의 심장부를 외부 라이브러리에서 손으로 복사해 온다 cp $CUDACONV/src/conv_util.cu $CUDACONV/src/filter_acts.cu $CUDACONV/src/weight_acts.cu $CUDACONV/src/img_acts.cu build/src cp $NVMATRIX/src/nvmatrix.cu $NVMATRIX/src/nvmatrix_kernels.cu $NVMATRIX/src/gpu_locking.cpp build/src cd build && make -j kepler=1 $* && cd ..
filter_acts.cu, weight_acts.cu, img_acts.cu. 순전파·가중치갱신·역전파를 담당하는 세 개의 손으로 깎은 커널. nvmatrix라는 자작 행렬 라이브러리. 그리고 저 gpu_locking.cpp ― 두 장의 카드가 서로의 발을 밟지 않도록 GPU에 자물쇠를 채우는 코드다. cuDNN? 엔비디아가 이런 걸 알아서 해주는 라이브러리는 2014년에야 나온다. 2012년의 다락방엔 그런 온실이 없었다. 라이브러리가 없으면, 라이브러리를 직접 짜는 수밖에.
빌드 설정(Makefile-distrib)을 보면 시대의 지문이 더 선명하다.
GENCODE_ARCH := -gencode=arch=compute_20,code="sm_20,compute_20"
compute_20. 페르미 세대(Compute Capability 2.0)를 겨냥해 컴파일하라는 명령이다. 요즘 카드가 compute_90을 넘나드는 걸 생각하면, 이건 흑백 필름으로 찍은 사진이다. 그리고 그 안에는, 논문 한 편이 학계를 뒤엎게 만든 바로 그 물건 ― ReLU ― 이 수줍은 C++ 클래스로 앉아 있다.
class ReluNeuron : public Neuron {
virtual void _activate() {
_inputs->apply(ReluOperator(), *_outputs); // 음수는 0으로, 그뿐
}
“음수는 0으로 만들어라.” 그 한 줄이, 수십 년간 신경망의 발목을 잡던 학습 정체를 풀어버린 열쇠였다. 위대한 발견은 대개 이렇게 시시하게 생겼다.
정작 나를 오래 붙든 건, 정답을 채점하는 커널이었다. src/layer_kernels.cu 안, kMultiSoftmaxCost_engs라는 함수에는 이런 주석이 달려 있다.
/*
* This routine uses energies to determine top-1 score because they're more
* accurate than top-n probabilities, which have numerical errors in them.
*/
__global__ void kMultiSoftmaxCost_engs(float* probs, float* energies, ...
float* top5Probs, ...) {
“확률값(top-n)에는 수치 오차가 끼니, 점수는 에너지값으로 매긴다.” 이미지넷 대회의 그 유명한 ’top-5 정확도’가, 부동소수점 반올림 오차와 싸우는 이런 진창 속에서 계산되고 있었던 것이다. 세상을 바꾼 15.3%라는 숫자 뒤에는, 소수점 아래에서 새는 오차를 틀어막던 이런 이름 없는 노동이 깔려 있다.
파이썬 쪽도 정겹다. convnet.py를 열면 파이썬 2 시절의 print가 괄호도 없이 튀어나오고, 카드 종류에 따라 다른 라이브러리를 불러오는 분기가 있다.
lib_name = "_ConvNet_k20x" if is_kepler_machine() else "_ConvNet" print "=========================" print "Importing %s C++ module" % lib_name
파이썬은 어디까지나 얇은 껍데기였다. 데이터를 밀어 넣고, 실험을 조립하고, C++ 엔진을 호출하는 관리자. 진짜 막노동은 그 아래 CUDA에서 벌어졌다. 실행 스크립트 run4.sh에 박제된 경로 ― /ais/gobi3/u/kriz/tmp/ConvNet__2012-08-03_14.28.23 ― 는 토론토의 ‘gobi’ 클러스터에서, 2012년 8월 3일 오후 2시 28분에 시작된 어느 학습 세션의 타임스탬프다. 그날 이후 에폭(epoch)은 22 → 49 → 66 → 73 → 81 → 95로, 학습률을 손으로 꺾어가며 늘어난다. 자동화된 하이퍼파라미터 튜닝? 없다. 사람이, 로그를 노려보다, 손으로 숫자를 바꿨다.
5. 48킬로바이트 위에서 6천만을 굴리는 법 ― 진짜 코드
앞에서 나는 “6천만 개의 파라미터를 48킬로바이트짜리 선반 위에서 한 바이트씩 굴렸다”고 썼다. 문학적 과장이라 여겼다면, 저장소가 나를 정정해 준다. 그것은 수사가 아니라 문자 그대로의 구현이었다.
먼저 오해 하나를 풀자. 6천만 파라미터가 48킬로바이트 안에 들어간 것이 아니다. 파라미터는 3기가바이트 전역 메모리(global memory)에 산다. 48킬로바이트는 각 연산 블록이 쓰는 초고속 ‘작업대’(shared memory)다. 문제는 이 작업대가 손바닥만 해서, 거대한 데이터를 한 번에 올릴 수 없다는 것. 그래서 코드는 데이터를 잘게 썰어 작업대에 올리고, 계산하고, 비우고, 다음 조각을 올리는 일을 지겹도록 되풀이한다. 이 반복(iteration)이야말로 GPU 프로그래밍의 알파이자 오메가다. 저장소는 그 반복을 세 겹으로 보여준다.
증거 ① ― 작업대를 32칸씩 채우는 타일 루프. src/layer_kernels.cu의 한 커널은 공유 메모리를 손수 할당하고, 데이터를 딱 32개씩 실어 나른다.
extern __shared__ float shmem[]; // 48KB 작업대를 직접 할당
...
for (int l = 0; l < numOut / 32; ++l) { // 32개씩 나눠 반복한다
__syncthreads(); // 코어들, 다 같이 멈춰 발 맞추기
// 작업대에 32개를 올린다
shE[(r + loadY) * (32 + 1) + loadX] = __expf(energies[r * numOut]);
printf("%fn", energies[r * numOut]); // (2012년의 디버깅 흔적, 지우지 못한 채 박제됨)
__syncthreads();
// ... 32행을 계산하고, 작업대를 비우고, 다음 32개로 ...
}
(32 + 1)이라는 어정쩡한 숫자에 눈이 멈춘다. 32가 아니라 33이다. 실수가 아니다. 여러 코어가 같은 메모리 은행에 한꺼번에 몰려 충돌(bank conflict)하는 것을 피하려고, 일부러 한 칸을 비워 어긋나게 깔아 놓은 것이다. 이런 ‘1칸짜리 미신’ 같은 최적화가, 3기가바이트 감옥 안에서 속도를 쥐어짜던 장인의 손끝이다. 그리고 저 printf ― 값을 화면에 찍어보던 디버깅 코드가 지워지지 않은 채 인류의 유산으로 공개됐다. 위대한 코드에도, 지우다 만 낙서가 있다.
증거 ② ― 6천만의 수정량을 조각내어 더하다. 파라미터를 얼마나 고칠지 계산한 ’기울기(gradient)’를 한꺼번에 담을 임시 공간조차 빠듯하다. 그래서 src/layer.cu의 합성곱 층은 partialSum이라는 손잡이로 이걸 조각낸다.
NVMatrix& tgt = _partialSum > 0 ? _weightGradTmp : _weights[inpIdx].getGrad();
convWeightActs(*_inputs[inpIdx], v, tgt, ..., _partialSum, scaleTargets, scaleWGrad);
if (_partialSum > 0) {
_weightGradTmp.reshape(_modules / _partialSum, ...); // partialSum 크기로 잘라
_weights[inpIdx].getGrad().addSum(_weightGradTmp, ...); // 조각들을 다시 합산한다
}
설정 파일을 다시 열면, 이 조각의 크기가 층마다 손으로 다르게 매겨져 있다. 얕은 첫 층은 partialSum=5, 다음은 partialSum=3, 깊은 층들은 partialSum=13. 왜 하필 5이고 3이고 13인가? 원리로 도출된 값이 아니다. 각 층의 크기와 3기가바이트 예산 사이에서, 밤새 돌려보며 ’안 터지는 숫자’를 손으로 찾아낸 것이다. 자동 튜닝은 없었다. 사람의 감(感)과 로그와 커피가 있었을 뿐이다.
증거 ③ ― 그리고 결국, 6천만을 한 줄로 훑는다. 모든 계산이 끝나면 6천만 파라미터를 갱신해야 한다. 스레드를 6천만 개나 띄울 수는 없으니, 유한한 코어가 성큼성큼 보폭을 넓혀 전체를 밟고 지나가는 ’격자 보폭 루프(grid-stride loop)’를 쓴다. src/weights.cu의 모멘텀 갱신 커널이다.
__global__ void superMomUpdate(float* V, float* g, ..., const int numElements) {
const int tidx = blockIdx.x * SUPERMOM_THREADS + threadIdx.x;
for (int t = tidx; t < numElements; t += gridDim.x * SUPERMOM_THREADS) {
V[t] = 0.9 * V[t] + eps * __fdividef(g[t], sqrtf(G2[t] - G[t] + superEps));
}
}
for (t = tidx; t < numElements; t += ...). 이 한 줄이, 6천만 개(numElements)의 파라미터를 유한한 코어가 보폭을 넓혀 남김없이 밟고 지나가는 방식이다. 0.9 * V[t]는 관성(모멘텀) ― 이전에 가던 방향을 90% 유지하며 미끄러지듯 내려가라는 명령이다.
이 세 겹의 반복이, 128장짜리 미니배치(convnet.py의 기본값 minibatch_size=128) 위에서 돌아간다. 120만 장을 128장씩 끊어 한 바퀴 도는 것이 1에폭(epoch)이고, 앞서 본 run4.sh의 22 → 49 → 66 → 73 → 81 → 95라는 숫자가 바로 그 바퀴 수다. 매 배치마다 작업대를 32칸씩 채우고(①), 기울기를 조각내 합치고(②), 6천만을 한 줄로 훑어 갱신한다(③). 이 지겨운 삼중 반복을 수십만 번 되풀이하고 나서야, 3기가바이트 감옥은 비로소 개와 고양이를 가르는 눈을 떴다.
정직하게 한마디 덧붙인다. 합성곱의 가장 안쪽 타일링 커널(convWeightActs 등)은 이 저장소가 아니라 크리제프스키의 별도 라이브러리 cuda-convnet2에서 빌드할 때 끌어온다(build.sh가 그 심부름을 한다). 그러나 저장소가 직접 품은 위 코드들만으로도, 48킬로바이트라는 작업대 위에서 6천만을 굴리던 그 규율은 조금의 과장도 없이 드러난다. 제미니가 지어낸 매끈한 가짜 커널과 달리, 진짜 사료는 이렇게 (32+1)의 미신과 지우다 만 printf까지 달고 있다. 사료의 값어치는 매끄러움이 아니라 흉터에 있다.
6. 지문 위에 덮인 도장 ― 왜 구글은 480억을 냈는가
발굴의 마지막에, 나는 두 개의 저작권 표시를 나란히 놓고 한참을 들여다봤다.
저장소 맨 위 LICENSE 파일의 첫 줄:
Copyright 2023 Google LLC.
그런데 그 안, src/의 모든 .cu 파일 머리에는 이렇게 적혀 있다:
Copyright (c) 2011, Alex Krizhevsky ([email protected])
2011년 한 대학원생의 지메일 주소 위에, 2023년 거대 기업의 법인명이 도장처럼 찍혀 있다. 이 두 줄의 간극이, 2012년 겨울 타호 호숫가의 그 유명한 경매를 설명한다. 힌튼과 두 제자가 세운 직원 셋짜리 유령 회사 DNNResearch는, 사무실도 사업계획도 없이 오직 이 코드 하나를 들고 나섰다. 매시간 100만 달러씩 호가가 오른 이메일 경매의 종착지는 약 4,400만 달러 ― 당시 환율로 480억 원. 구글이 480억을 낸 것은 ’강아지 사진 분류기’가 아니었다. 저 160개의 실패한 설정 파일과, 손으로 깎은 CUDA 커널과, 3기가바이트 안에 뇌를 반으로 접어 넣은 그 지독한 노동 자체였다. 창의성은 논문의 수식이 아니라, 수식을 3,072메가바이트짜리 감옥 안에서 기어이 작동시킨 그 진창 속에 있었기 때문이다.
인공지능은 두 번 여름을 맞았다가 두 번 겨울에 얼어 죽었다. 1950년대의 첫 여름도, 1980년대의 두 번째 여름도, 화려한 약속을 남기고 시들었다. 2012년의 이 투박한 코드가 연 것은 세 번째 여름이었다. 그리고 이번 여름은, 앞선 둘과 달리 겨울로 돌아가지 않았다. 성능이 안 나와 수십 년을 허덕이던 신경망이, 하필 게임용 칩이라는 값싼 삽 한 자루를 만나 마침내 영원한 개화(開花)로 넘어간 것이다. 그 변곡점의 지질학적 증거가, 지금 누구나 내려받을 수 있는 이 45,042줄의 설정 파일과 8,449줄의 CUDA 안에 고스란히 화석으로 남아 있다.
최초로 융합하는 자는 늘 맨땅에서 개고생한다. 서로 만날 이유가 없던 두 세계를 처음 포개는 일이란 원래 이렇게, 오래 걸리고 지저분하고 자주 망하도록 되어 있다. 그러니 정책이 할 일은 그 개고생을 ’표준 미달’이라 결재판에서 탈락시키는 것이 아니라, 160번을 망해도 161번째를 시도할 수 있는 값싼 삽과 너른 모래밭을 깔아주는 것이다. 나는 오늘 그 증거를, 남의 요약이 아니라 내 손으로 판 지층에서 확인했다.
기계는 사료를 지어낼 수 있다. 그러나 지어낸 사료 위에는 아무것도 세울 수 없다. 역사학자의 첫 번째 의무는, 삽을 들고 직접 내려가 보는 것이다.
(인용한 모든 코드는 computerhistory/AlexNet-Source-Code 저장소의 실제 파일에서 발췌했다. 가공하지 않았다.)