들어가며
요즘 AI 때문에 GPU 이야기를 안 들어본 개발자가 없을 거예요. “GPU로 학습 돌린다”, “CUDA 커널을 짠다” 같은 말은 많이 들었는데, 정작 우리가 GPU한테 작업을 시키면 그 안에서 실제로 무슨 일이 일어나는지는 의외로 잘 모르고 넘어가거든요. 오늘은 그 ‘GPU 커널 실행’의 내부를 한 단계씩 따라가 볼게요.
먼저 용어부터. ‘커널(kernel)’이라고 하면 운영체제 커널을 떠올리기 쉬운데, GPU에서 말하는 커널은 전혀 달라요. GPU에서 커널은 그냥 ‘GPU 위에서 돌아가는 함수 하나’를 뜻해요. CUDA는 엔비디아 GPU에 이런 함수를 돌리게 해주는 도구이자 언어고요.
코드 한 줄이 실행되면
C++ 코드에서 보통 이렇게 커널을 불러요:
myKernel<<<1024, 256>>>(data);
이 한 줄 안에 어마어마한 일이 숨어 있어요. 꺾쇠 안의 숫자(1024, 256)는 “이 함수를 몇 갈래로 쪼개서 동시에 돌릴까”를 정하는 거예요. 여기선 256개짜리 묶음(블록)을 1024개 만들라는 뜻이니, 무려 26만 개가 넘는 작은 일꾼(스레드)이 같은 함수를 동시에 실행하게 돼요. CPU가 한 번에 몇 개 안 되는 일을 빠르게 처리하는 ‘소수 정예’라면, GPU는 하나하나는 느려도 어마어마하게 많은 일꾼이 한꺼번에 달려드는 ‘인해전술’인 셈이죠.
이 한 줄을 만나면 먼저 ‘CUDA 런타임’이라는 중간 관리자가 나서요. 이 친구가 “어떤 함수를, 일꾼을 몇 명으로, 어떤 데이터로 돌려라”라는 주문서를 만들어요. 그런데 여기서 중요한 사실 하나. 이 주문은 곧바로 실행되는 게 아니라 GPU 앞에 줄을 서요(이 줄을 ‘스트림’이라고 불러요). CPU는 주문서만 GPU 쪽 줄에 던져놓고, 결과를 기다리지 않고 바로 자기 다음 할 일을 해요. 이게 ‘비동기(asynchronous)’ 실행이에요. 카페에서 주문만 하고 자리에 앉아 다른 일 하는 것처럼요.
드라이버와 하드웨어로 내려가면
주문서는 ‘GPU 드라이버’를 거쳐요. 드라이버는 운영체제와 GPU 하드웨어 사이의 통역사예요. 이 통역사가 주문을 GPU가 알아듣는 명령어로 바꿔서, PCIe라는 고속도로(메인보드와 GPU를 잇는 통로)를 통해 GPU 안의 명령 대기열로 보내요.
GPU 안으로 들어가면 이제 진짜 일이 시작돼요. GPU에는 SM(스트리밍 멀티프로세서)이라는 작업장이 수십 개 들어 있어요. GPU 안의 스케줄러가 우리가 만든 1024개의 블록을 이 작업장들에 나눠서 배치해요. 한 작업장이 동시에 여러 블록을 맡기도 하고요.
그리고 핵심 개념이 하나 더 있어요. ‘워프(warp)’예요. GPU는 스레드를 한 개씩 다루지 않고 32개씩 한 묶음으로 다뤄요. 이 32개 묶음이 워프인데, 같은 워프 안의 32개 일꾼은 ‘완전히 똑같은 명령’을 동시에 실행해요. 들고 있는 데이터만 각자 다르고요. 이걸 어려운 말로 SIMT라고 부르는데, 쉽게 말하면 “구령 하나에 32명이 똑같은 동작을 하는 제식훈련” 같은 거예요. 이 방식 덕분에 GPU가 엄청난 병렬 처리를 효율적으로 해내는 거죠.
CPU 방식과 뭐가 다르냐면
여기서 흥미로운 함정이 하나 있어요. 같은 워프 안의 일꾼들이 if 문에서 서로 다른 길로 갈라지면(이걸 ‘워프 다이버전스’라고 해요), GPU는 양쪽 길을 차례로 다 거쳐야 해서 속도가 뚝 떨어져요. 32명이 똑같이 움직여야 빠른데, 절반은 왼쪽 절반은 오른쪽으로 가라고 하면 비효율이 생기는 거죠. 그래서 GPU 코드를 잘 짠다는 건, 가능한 한 모든 일꾼이 같은 길을 가게 만드는 거예요. CPU 코드 최적화와는 사고방식이 완전히 다르죠.
한국 개발자에게 주는 시사점
AI 모델을 학습시키든, 대량 데이터를 처리하든, 요즘은 GPU를 안 건드릴 수가 없어요. 그런데 많은 분들이 파이토치(PyTorch) 같은 라이브러리 뒤에 가려진 이 내부 동작을 모른 채 “왜 이렇게 느리지?”만 반복하거든요. 비동기로 줄을 선다는 사실 하나만 알아도, 왜 측정한 시간이 이상하게 나오는지(GPU가 아직 일을 끝내지도 않았는데 CPU가 먼저 시간을 재버려서) 이해가 되고요. 워프 개념을 알면 왜 특정 코드가 갑자기 느려지는지도 보여요.
당장 CUDA 커널을 직접 짤 일이 없더라도, 이 흐름을 이해해두면 GPU 성능 디버깅이나 비용 최적화에서 한 차원 깊은 시야를 갖게 돼요.
정리하면
GPU 커널 실행은 ‘주문서 작성 → 줄 서기 → 드라이버 통역 → 작업장 배치 → 32명 단위 동시 실행’이라는 여정을 거쳐요. 여러분은 GPU를 ‘마법 상자’로 두고 계신가요, 아니면 그 안을 들여다볼 준비가 되셨나요?
🔗 출처: Hacker News