스탠퍼드 CS149 병렬 컴퓨팅 2023 제7강: GPU 아키텍처와 CUDA 프로그래밍

⏱ 02:37GPU는 원래 장면의 기하 구조·광원·카메라 정보를 받아 영상을 만드는 그래픽 칩이었다. 수백만 픽셀의 색을 초당 60회 이상 계산해야 하므로 멀티코어와 SIMD가 발전했다. 단일 스레드 CPU의 클럭·전력·명령어 수준 병렬성이 한계에 이른 반면, GPU는 늘어난 트랜지스터를 코어 확대로 활용했다. 연구자들은 화면을 삼각형 두 개로 덮어 픽셀 프로그램을 범용 계산에 이용했고, 브룩을 거쳐 CUDA처럼 직접 병렬 작업을 요청하는 인터페이스로 발전했다.

⏱ 15:38CUDA는 동일한 프로그램을 서로 다른 ID로 여러 번 실행하는 SPMD 모델이다. CUDA 스레드는 하드웨어 실행 문맥과 구분해야 하며, 스레드 블록은 ISPC 태스크와 비슷한 계층을 이룬다. 12×6 행렬을 4×3 블록으로 처리하면 스레드 72개가 필요하다. 11×5로 바꾸고 블록 크기를 유지하면 올림 때문에 불필요한 스레드도 생기므로, 각 스레드가 인덱스를 계산하고 경계를 검사해야 한다. 배열 원소마다 자동으로 실행되는 모델이 아니라 프로그래머가 스레드를 일괄 생성하는 모델이라는 점이 핵심이다.

⏱ 26:01기본 메모리 모델에서는 CPU와 GPU의 주소 공간이 분리된다. 호스트 배열과 cudaMalloc으로 만든 디바이스 배열 사이를 cudaMemcpy로 옮기고, 커널에는 디바이스 포인터를 전달한다. 별도 GPU 카드에서는 서로 다른 DRAM 사이의 PCIe 전송이므로 비용이 크며 비동기 복사로 지연을 숨길 수 있다. GPU 전역 메모리는 여러 스레드가 접근하지만, 지역 변수는 스레드별이고 공유 변수는 블록별이다. 이 구분은 협업할 스레드를 같은 위치에 배치하여 지역성을 확보하는 수단이기도 하다.

⏱ 36:211차원 합성곱에서는 이웃 스레드가 읽는 입력이 겹친다. 출력 128개에 필요한 입력 130개를 블록의 공유 배열로 한 번 가져오면 중복 전역 접근을 줄일 수 있다. 스레드 128개가 하나씩 적재하고, 처음 두 스레드는 나머지 두 원소도 맡는다. 이후 __syncthreads()로 모든 적재 완료를 확인한 뒤 각자 출력을 계산한다. 이 장벽을 빼면 이웃이 아직 쓰지 않은 값을 읽을 수 있다. 재사용은 한 스레드 내부가 아니라 이웃 스레드 사이에서 생기며, 연속 주소 접근과 데이터 배치도 성능에 영향을 준다.

⏱ 53:12CPU에서는 컴파일러가 SIMD 명령어를 만들지만 GPU는 같은 프로그램 위치에 있는 CUDA 스레드들을 하드웨어가 묶어 암묵적 SIMD로 실행한다. NVIDIA의 워프는 스레드 32개분의 실행 문맥이다. PC가 다른 스레드는 마스킹하며, 16개 ALU로 32개 스레드를 처리하는 예에서는 두 클럭을 쓴다. 볼타의 SM 예시는 네 서브코어와 64워프, 즉 실행 문맥 2,048개를 갖고 클럭마다 최대 네 워프를 진행시킨다. 많은 문맥을 유지하는 목적은 지연 숨기기이며, 모두가 같은 순간에 연산한다는 뜻은 아니다. V100 예시에는 SM 80개가 있고, SM당 64워프와 워프당 32스레드를 곱하면 약 163,000개 문맥을 유지한다. 강의의 수치는 해당 구조를 설명하는 예이며, 블록 크기와 저장 공간 요구를 직접 지정하면 하드웨어에 대한 가정도 생긴다.

⏱ 47:23커널 실행 요청은 수많은 블록을 하드웨어 스케줄러에 맡기는 작업 큐와 비슷하다. 약 100만 스레드와 8K 블록을 요청해도 그만큼의 하드웨어 문맥이 한꺼번에 필요한 것은 아니다. 컴파일된 코드의 스레드 수·공유 메모리 요구량을 보고, 자원이 있는 SM에 블록을 올리고 완료 후 다음 블록을 배치한다. 실행 문맥이 남아도 공유 메모리가 부족하면 기다려야 한다. 같은 커널의 블록은 요구량이 균일하므로 자원 재사용이 단순하고, 코어 수가 다른 GPU에도 같은 프로그램을 배치할 수 있다.

⏱ 1:09:49블록은 동기화와 지역성의 단위이므로 같은 SM에서 구성 스레드들이 함께 살아 있어야 한다. 실행 문맥 128개인 코어에 256스레드 블록을 절반씩 실행하면, 첫 절반이 장벽에서 문맥을 차지한 채 기다려 나머지가 실행되지 못한다. 블록 간에는 전역 atomicAdd로 히스토그램 빈을 갱신할 수 있지만 실행 순서는 보장되지 않는다. 한 블록이 다른 블록의 갱신을 기다리는 코드는 대기 블록이 먼저 상주할 때 교착될 수 있다. 학습의 핵심은 SPMD 모델과 SIMD 구현을 구분하고, 자원 배치·지역성·동기화 조건을 함께 추적하는 것이다.

⏱ 00:06시작하겠습니다. 다른 일들은 잘되고 있나요? 이제 4주 차이고, 오늘은 GPU를 이야기하겠습니다. 먼저 현대 GPU 컴퓨팅에 이르게 된 역사를 조금 살펴봅니다. NVIDIA와 AMD가 주로 퀘이크(Quake) 같은 게임을 하도록 만들던 칩이 어떻게 NVIDIA를 인텔보다 훨씬 큰 시가총액 1조 달러 기업으로 만들고, 모두가 이 장치를 사려고 경쟁하는 상황으로 이어졌을까요? 그 역사를 살펴본 다음 NVIDIA가 제공하는 언어로 GPU를 프로그래밍하는 방법을 다룹니다. 요즘 GPU 코드는 대부분 CUDA로 작성합니다.

⏱ 01:01CUDA는 ISPC와 아주 비슷하게 느껴질 것입니다. ISPC의 역사를 다룬 블로그 글을 읽어 보면, ISPC는 사실 CUDA에 대한 반응으로 등장했습니다. 사람들이 GPU를 CUDA로 프로그래밍하는데 CPU도 같은 방식으로 프로그래밍하면 안 되겠느냐고 생각해서 ISPC를 만든 것입니다. 따라서 프로그래밍 모델은 익숙할 것입니다. 이어서 현대 GPU의 구조와 그 구조가 CUDA 코드를 실행하는 방식을 살펴봅니다. 첫째·둘째 주 수업과 멀티스레딩, SIMD, 멀티코어 실행을 상세히 다룬 필기 과제, CPU에서 코드를 실행했던 경험을 떠올리세요. 오늘 새로운 개념은 없습니다. 같은 아이디어를 현대 GPU에서 더 큰 규모로 구현할 뿐입니다.

⏱ 02:15여러분 중 CS 148을 들었거나 이 수업 뒤에 CS 148을 듣는 사람, 또는 저와 CS 248을 들었던 사람이 있나요? 일부는 그래픽스와 영상이 렌더링되는 방식을 조금 알 것입니다. 나머지 여러분을 위해 CS 248을 슬라이드 다섯 장 정도로 설명하겠습니다. 원래 GPU 설계의 목적은 하나의 특정 문제를 푸는 것이었습니다. 표면의 기하 구조, 광원의 위치, 가상 카메라의 위치처럼 장면을 수학적으로 기술하면, 빛이 장면 안에서 반사되는 과정과 여러 재질을 시뮬레이션하여 실제 카메라가 그곳에 있었다면 보았을 영상을 만드는 것입니다. 그래픽 칩의 목표였고, 실제로 아주 잘하는 일입니다.

⏱ 03:23여기 보이는 영상은 이제 꽤 오래된 2015년 영상입니다. 당시 고급 GPU로도 초당 60프레임을 쉽게 렌더링할 수 있었습니다. 좀 더 현대적인 그래픽스로 넘어가면, 이 예시는 아주 고급 GPU에서 실시간으로 렌더링할 수 있는 에픽(Epic)의 나나이트(Nanite) 데모입니다. 이런 결과를 계산하는 소프트웨어를 작성해 보라고 한다면, 실제 연산 작업 자체는 비교적 간단하게 설명할 수 있습니다. CS 248을 약 45초로 압축해 보겠습니다. 핵심은 여기 보이는 것과 같은 표면을 렌더링하는 일입니다.

⏱ 04:07예전의 GPU 사용 방식, 또는 그래픽스 작업 방식은 GPU에 표면의 표현을 제공하는 것입니다. 컴퓨터 그래픽스에서는 보통 삼각형 메시(triangle mesh)로 표면을 나타냅니다. 모든 삼각형 꼭짓점의 3차원 좌표 목록과 카메라의 위치·방향을 GPU에 주고, CS 248에서 배우는 수학을 적용해 각 꼭짓점이 화면 어디에 투영되는지 계산하게 합니다. 여기서 카메라를 들면 각 꼭짓점이 화면의 어디에 나타날지를 구하는 것입니다. 삼각형의 화면상 위치를 알게 되면, 그 삼각형들이 덮는 영상의 모든 픽셀에 대해 색을 계산해야 합니다.

⏱ 05:05삼각형의 색을 계산하려면 이 방의 표면들을 생각해 보세요. 색은 그 표면을 이루는 재질에서 나옵니다. 의자, 대나무 무늬의 표면, 반짝이는 금속 문손잡이에서 빛이 어떻게 반사되는지를 시뮬레이션하다 보면 세상에 정말 다양한 재질이 있음을 깨닫습니다. 여기 예시들은 서로 비슷해 보이지만 재질 표현은 훨씬 정교해질 수 있습니다. 신발의 반짝임이나 사람 피부를 보세요. 빛은 피부 안으로 들어가 여러 번 산란한 뒤 다른 위치로 나옵니다. 재질마다 특성이 너무 다르므로 단순히 빨강이나 파랑이라고 지정하는 대신, 이런 빛의 작용을 계산하는 프로그램으로 삼각형의 색을 구하고 싶었던 것입니다.

⏱ 06:082000년대 초에는 화면의 픽셀마다 실행할 작은 프로그래밍 언어들이 개발되기 시작했습니다. 프로그램은 현재 재질과 표면의 방향에 관한 정보를 받습니다. 코드의 세부 내용은 여기서 중요하지 않지만, 표면의 어느 위치인지를 나타내는 2차원 좌표, 표면 바깥을 향하는 벡터인 법선(normal), 텍스처 맵 정보 등이 들어 있습니다. 표면 방향과 의자 표면에 씌우고 싶은 영상이 주어지면, 알고리즘이 약간의 연산을 수행해 픽셀의 색을 구합니다. 출력으로 반환되는 벡터가 색입니다. Kd는 RGB를 담은 vec3이고, 마지막 성분은 알파(alpha)입니다.

⏱ 07:10이 함수를 모든 원소, 즉 모든 픽셀에 실행하면 여기와 같은 영상이 나옵니다. 수업에서 다룬 기본적인 데이터 병렬 연산입니다. 픽셀별 입력을 마련하고 서로 완전히 독립적인 함수를 각 픽셀에 실행하여 색을 계산합니다. 이것이 CS 248의 핵심입니다. GPU에 멀티코어와 SIMD가 들어간 이유도 여기에 있습니다. 4K 영상에는 수백만 픽셀이 있고, 이를 초당 60회 이상 처리해야 합니다. 초당 수억 개에서 수십억 개 픽셀을 처리하려면 색을 더 빨리 계산할 코어와 ALU가 필요했으므로 그 수를 계속 늘렸습니다.

⏱ 08:06거의 정확히 20년 전으로 돌아가 첫 강의를 떠올려 보세요. 녹색 선은 칩당 트랜지스터 수가 계속 늘어나는 모습을 나타냈지만, 트랜지스터를 더 빠른 단일 스레드 프로세서로 바꾸던 방법들은 한계에 이르렀습니다. 4GHz 프로세서의 전력 소모가 치솟아 클럭을 계속 높일 수 없었고, 스레드에서 더 찾을 명령어 수준 병렬성(ILP)이 없어 슈퍼스칼라도 계속 확대할 수 없었습니다. 반면 GPU는 이미 병렬화를 택했으므로 녹색 선이 올라갈수록 코어를 더 넣으면 되었습니다. 스탠퍼드, 노스캐롤라이나 대학교 등의 연구자들은 인텔 프로세서는 빨라지지 않는데 GPU는 트랜지스터 수가 두 배가 될 때마다 코어가 늘고, 이미 색을 계산하는 작은 프로그램도 실행한다는 점에 주목했습니다.

⏱ 09:11그래서 사람들이 우회 기법을 쓰기 시작했습니다. 512×512 화면에 영상을 렌더링한다고 합시다. 삼각형으로 화면 전체를 덮고, 색을 구하던 프로그램에 대신 유용한 일을 시킵니다. 물리 시뮬레이션에서 입자를 한 시간 단계 진행시키거나 단백질 접힘을 계산하는 식입니다. 삼각형 두 개를 그리는 행위가 사실상 512²번의 함수 호출을 만드는 셈입니다. 표면 색을 계산하라고 마련된 함수 안에 실제로 원하는 작업을 넣는 것입니다. 약 20년 전 이런 작업을 하던 사람들은 RGBA 출력을 유체 시뮬레이션의 XYZ 위치로 해석할 수도 있다고 생각했습니다. GPU가 대단히 빨라 C++의 스레드 코드보다 훨씬 빠르게 실행될 수 있다는 사실도 확인했습니다.

⏱ 10:242004년 스탠퍼드의 한 연구 프로젝트는 이 방식이 너무 어색한 우회라고 지적했습니다. 어차피 GPU를 병렬 프로세서로 쓰고 데이터 병렬성도 알고 있다면, 그에 맞는 제대로 된 프로그래밍 시스템을 만들 수 없겠느냐는 것입니다. 이 프로젝트는 목요일에 더 자세히 다룰 데이터 병렬 프로그래밍 모델을 사용했습니다. 예를 들어 A와 B를 입력받아 A를 일정량 배율 조정하는 scale 함수가 있다고 합시다. 이를 스칼라 값이 아니라 원소들의 집합에 적용합니다. 여기서 보여 주는 구문은 설명을 위한 것으로, 숫자 1,000개가 담긴 벡터의 모든 원소에 scale을 적용한다고 생각하면 됩니다. 이를 스트림 프로그래밍 모델 또는 데이터 병렬 프로그래밍 모델이라고 부릅니다.

⏱ 11:23이 코드는 브룩(Brook)이라는 언어에서 실제로 유효한 코드였습니다. 소스 대 소스 컴파일을 거쳐 삼각형 두 개를 그리고 GPU를 이용하는 그래픽스 프로그램으로 바뀌었습니다. 사용자 입장에서는 우회 기법을 덜 의식해도 되었지만, 구현 자체는 여전히 완전한 우회였습니다. 당시 NVIDIA는 프로세서를 계속 빠르게 만들면서, 칩이 코드를 실행할 수 있으니 그래픽스 이외의 일도 할 수 있지 않을까 생각하기 시작했습니다. 이런 프로젝트 등에서 영감을 얻어, 이미 칩이 실행할 수 있는 범용 코드를 사람들이 직접 작성할 수 있도록 NVIDIA 스스로 지원하기로 한 것입니다.

⏱ 12:07CPU의 작동 방식을 떠올려 봅시다. 코어가 두 개이고 실행 문맥도 총 두 개인 멀티코어 CPU가 있다고 합시다. 이 프로세서에서 코드를 실행하려면 어떻게 할까요? 다시 말해 운영체제는 무엇을 할까요? 아직 운영체제 수업을 듣지 않은 사람이 많겠지만, 그 수업에서 운영체제를 만든다고 가정해 보세요. 두 코어를 가진 이 컴퓨터에서 서로 다른 프로그램 두 개를 실행하려면 어떻게 시작해야 할까요? 프로그램 두 개를 컴파일했으므로 명령어 흐름도 두 개 있습니다. 운영체제로서 무엇을 해야 할까요? 한 학생이 각 명령어 흐름을 해석해서 각 코어에 배정하면 되겠다고 답합니다.

⏱ 13:05맞습니다. 운영체제는 첫 명령어 흐름을 주고 코어의 레지스터 값을 초기화한 뒤, 다음에 실행할 명령어를 프로그램의 첫 위치로 지정하고 실행시킵니다. 그렇게 스레드가 실행됩니다. 다른 프로그램의 또 다른 스레드도 예를 들어 코어 0에 초기 상태를 넣고 실행할 수 있습니다. 프로그램 바이너리를 주고 스레드 상태를 초기화한 다음 시작하는 일을 스레드 단위로 하는 것입니다. 이것이 운영체제가 CPU 코드를 실행하는 방식입니다. 하지만 NVIDIA가 직접적인 GPU 접근을 제공하기 전에는 GPU에 스레드를 실행하라는 명령이 없었습니다. 문자 그대로 삼각형들을 그리고, 어떤 픽셀에 색을 칠해야 하는지 구한 다음 지정한 프로그램으로 그 픽셀의 색을 계산하라고 명령했습니다.

⏱ 14:04이 GPU 인터페이스가 그래픽스 파이프라인입니다. CS 248에서 다루는 알고리즘들이 그 안에 들어 있습니다. 소프트웨어가 삼각형 목록, 표면 색을 구하는 프로그램, 카메라 위치를 주면 파이프라인이 동작하면서 픽셀 색을 구할 때마다 그 프로그램을 실행합니다. CPU와는 계산을 시작하는 사고방식이 매우 다릅니다. 2007년 NVIDIA는 이 두 방식의 중간에 있는 인터페이스가 필요하다고 판단했습니다. 그래픽스와 관계없는 일을 하려고 삼각형을 그리는 것은 어색하지만, 사용자가 제각각 스레드를 만들게 하는 것도 바람직하지 않았습니다. 폭 32의 SIMD를 쓰고 단일 스레드 성능이 좋지 않은 GPU에서는 그런 임의의 스레드들이 잘 실행되지 않기 때문입니다.

⏱ 15:01그래서 지금도 존재하는 컴퓨트 모드(compute mode)라는 추상화를 만들었습니다. 하드웨어 인터페이스는 이렇습니다. foo 같은 작은 함수 프로그램을 작성하고 그것을 N개 실행해 달라고 요청합니다. 여기서는 myKernel이라고 불렀습니다. 데이터 병렬 문맥에서는 실행할 코드를 흔히 커널(kernel)이라고 합니다. GPU 칩에 커널 N개를 실행하라고 하고, 구체적인 실행 방법은 GPU가 알아서 결정하게 합니다. 어떤 프로그래밍 모델일까요? 전형적인 SPMD입니다. 어딘가에 스레드 ID가 들어 있는 동일한 코드를 주고 N개 사본을 시작해 실행하는 것입니다. 지금까지 본 프로그래밍 모델과 비슷하다고 생각하면 됩니다.

⏱ 15:51ISPC 같은 모델과 비슷하다는 뜻입니다. 그 모델을 구현한 언어가 CUDA입니다. 현대 GPU 프로그래밍은 기본적으로 C·C++과 같지만 스레드를 하나씩 만들지 않습니다. 프로그램을 주고 N개 사본을 실행해 달라고 하면 GPU가 그 사본들을 병렬화하는 방법을 결정합니다. 학생이 앞서 삼각형 두 개를 계산하던 방식과 어떻게 다른지, 이전에는 SPMD가 아니었는지 묻습니다. 이전에는 정말 OpenGL을 열어 삼각형을 특정 위치에 그리라고 했고, 모든 픽셀이 덮이도록 위치를 신중하게 정해야 했습니다. 그것은 우회 기법이었습니다. 이제는 제대로 된 SPMD 추상화로 사실상 같은 일을 하는 것입니다.

⏱ 16:53이것이 현재 CUDA에 이르기까지의 역사입니다. 이제 프로그래밍 모델을 설명하는 동안 이것을 데이터 병렬로 볼지 SPMD로 볼지, 메시지 전달 방식인지 공유 주소 공간 방식인지 생각해 보세요. ISPC와 이미 아는 개념들에 어떻게 대응시킬 수 있을까요? 오늘은 CUDA 용어를 쓰겠습니다. ISPC에는 프로그램 인스턴스의 갱(gang)이 있었습니다. 각 인스턴스를 논리적으로 스레드처럼 생각했지만, 실제 구현은 하드웨어 실행 문맥이 아니라 벡터 레인이었습니다. ISPC의 프로그램 인스턴스를 CUDA에서는 CUDA 스레드라고 부릅니다. 제가 명시적으로 하드웨어 실행 문맥이라고 말하지 않는 한, 오늘 강의의 스레드는 NVIDIA가 정의하는 CUDA 스레드를 뜻합니다.

⏱ 18:10ISPC 함수를 호출하면 프로그램 인스턴스 N개가 든 갱이 시작되고, N은 대개 컴파일 시점에 정해지며 흔히 8 정도였습니다. ISPC 태스크에서는 특정 함수나 코드를 N번 실행한다고 요청했습니다. CUDA도 비슷하게 생각하면 됩니다. 여기의 일반 C++ 호출은 배열 A, B, C를 인수로 CUDA의 matrixAdd 함수를 실행하라고 요청합니다. 배열은 보통의 C 스타일 배열입니다. 특이한 괄호 구문은 단일 N 대신 다차원으로 인스턴스 수를 지정합니다. 그래픽스와 텐서 연산에서 편리하기 때문입니다. 예를 들어 스레드를 4×3 격자로 구성할 수 있고, 스레드 ID도 다차원이 됩니다. 여기에 인스턴스들을 블록으로 묶는 구조가 하나 더 있습니다.

⏱ 19:39이 코드는 블록 하나를 4×3, 즉 스레드 12개로 구성하고, 블록 수를 한 방향으로 12/4개, 다른 방향으로 6/3개 만듭니다. 각각 12×6인 두 행렬을 원소별로 더하면서 CUDA 스레드 하나가 원소 하나를 맡도록 하려면 총 72개 스레드가 필요합니다. 그냥 72개를 만들라고 하는 대신 여기서는 4×3 블록으로 묶었습니다. 지금 스레드 블록을 소개하는 이유는 나중에 다시 쓰이기 때문입니다. 실제로 단순한 배열 두 개를 더하기만 한다면 1차원 스레드 블록을 골라 크기가 72인 블록 하나를 사용할 수도 있습니다.

⏱ 20:51기본 CUDA 구문을 보겠습니다. 앞 슬라이드의 호출 코드는 main.cpp에 있고, 지정한 수의 스레드 블록으로 matrixAdd를 호출합니다. CUDA의 matrixAdd 함수는 ISPC의 SPMD 프로그램과 매우 비슷합니다. programCount와 programIndex 대신 blockIdx, blockDim, threadIdx라는 내장 변수를 써서 자신만의 ID를 구합니다. threadIdx는 (0,0), (1,0)처럼 블록 내부의 스레드 위치이고, blockIdx는 블록을 나타냅니다. 전체에서의 위치를 구하려면 블록 크기도 알아야 합니다. 블록 안 스레드 ID는 ISPC의 프로그램 인스턴스 ID, 블록 ID는 태스크 ID와 비슷합니다. ISPC가 태스크와 그 안의 갱으로 구성되는 반면 CUDA는 여러 스레드 블록을 한꺼번에 만듭니다. 익숙한 개념에 다른 이름과 약간 다른 세부 사항이 붙은 것입니다.

⏱ 22:27matrixAdd에서는 각 스레드가 자신의 스레드 ID와 블록 ID로 입력 배열의 어느 원소를 맡는지 먼저 계산한 뒤 그 일을 수행합니다. 이 프로그램에서는 함수 인스턴스 72개가 동시에 실행되고, 행렬 원소마다 스레드 하나가 배정되어 전체 배열의 덧셈이 이루어집니다. 학생이 왜 2차원 방식으로 넘어왔는지 묻습니다. 실제로 이런 방식이 있으니 알려 주려는 것이며, 이 입문 예제들은 스레드·블록 ID를 모두 1차원으로 해도 잘 동작합니다. 다만 영상 처리, 텐서 처리, 그래픽스에서는 2차원 ID로 i와 j를 바로 얻는 것이 편리합니다. 1차원 ID로 N차원 텐서의 위치를 복원하려면 나눗셈을 많이 해야 하므로, 다차원 위치를 직접 주면 나눗셈을 줄여 메모리 주소를 효율적으로 계산할 수 있습니다.

⏱ 23:572차원 인덱스 자체가 근본적인 요구는 아닙니다. 해마다 처음에는 1차원 ID만 소개할까 고민하지만 결국 나중에 2차원을 또 설명해야 하므로 시간을 아끼려고 지금 함께 다룹니다. 이 시점에서는 두 세계를 떠올리세요. void main 같은 일반 C++ 코드는 CPU 스레드에서 실행되고, 여러 matrixAdd 인스턴스는 다른 세계에서 실행됩니다. ISPC에서는 갱을 같은 스레드 안의 SIMD 명령어로 구현한다고 했습니다. CUDA에서는 아직 구현 세부 사항을 깊이 다루지 않지만, CUDA 스레드들을 실제로 GPU에서 실행합니다. matrixAdd를 시작하는 호출이 내부 구현상 GPU 실행을 시작하라는 통신이 이루어지는 지점입니다.

⏱ 25:13학생이 이전 그림에서 블록 차원을 짚어 달라고 요청합니다. 이 예제에서는 꼭 필요하지 않지만, 72개 스레드를 그냥 실행하는 대신 일정한 수의 블록으로 나누었습니다. 각 블록은 4×3 스레드 격자이므로 스레드 ID가 2차원이고, 그림의 숫자들이 그 ID입니다. 이 슬라이드에는 블록 크기가 직접 표시되어 있지 않지만 blockDim은 4×3입니다. blockIdx는 현재 블록이 어느 블록인지를 나타냅니다. 결국 다차원 주소 지정이라고 생각하면 됩니다. 이해되나요?

⏱ 26:01CUDA 메모리 모델에서는 CPU에서 실행되는 void main에 CPU 자체의 주소 공간이 있습니다. 단순한 CUDA 모델에서는 디바이스, 즉 GPU도 CUDA 스레드가 접근하는 별도의 메모리 주소 공간을 가집니다. 따라서 C에서 배열을 할당하고 그 포인터를 CUDA에서 역참조하면 작동하지 않습니다. 주소 공간이 다르기 때문입니다. 예제의 CPU 쪽 void main은 보통의 C 할당으로 CPU 주소 공간에 배열을 만듭니다. 이어 cudaMalloc으로 GPU 쪽 디바이스 주소 공간에 배열을 할당하고, 라이브러리 호출로 CPU 공간의 값을 GPU 공간으로 옮깁니다. CUDA 커널에는 deviceA 포인터를 넘깁니다. GPU에서 실행되는 커널은 GPU 주소 공간의 메모리에 접근해야 하기 때문입니다.

⏱ 27:19이것이 CUDA의 기본적인 실무 구성입니다. 요즘 시스템에서는 C 포인터를 CUDA 커널에 바로 넘기기도 쉬워졌지만, 그 경우 GPU의 메모리 접근이 PCIe 같은 경로를 거쳐야 할 수도 있습니다. 지금은 명확하게 분리된 주소 공간 모델을 사용하겠습니다. 학생이 앞으로 할 계산을 알 때 CPU 데이터를 GPU로 미리 가져오는 것이 유용한 경우가 있는지 묻습니다. 질문 일부는 들리지 않지만 답은 그렇다는 것입니다. 여기 cudaMemcpy는 CPU 메모리에서 데이터를 옮깁니다. RTX GPU나 A100 같은 별도 GPU 카드가 있는 시스템에서는 실제 바이트들이 PCI 버스를 지나 GPU DRAM으로 이동합니다. 서로 다른 DRAM 두 개 사이의 이동이므로 느린 복사입니다.

⏱ 28:33이 복사를 비동기로 수행하는 방법도 있습니다. 지난 시간에 다룬 메시지 전달을 떠올려 보세요. 어떤 의미에서는 이 memcpy가 한 주소 공간에서 다른 주소 공간으로 보내는 메시지입니다. 따라서 지연 시간을 숨기려고 비동기로 수행하고 싶을 수 있습니다. 학생이 matrixAdd 예제로 돌아가 함수에 A, B, C를 넘기는 점을 묻습니다. 그 예제의 A, B, C는 모두 CUDA 디바이스에 할당한 배열입니다. 이 예제로 돌아와 주어서 좋습니다. 같은 예제를 조금 바꾸어 여러분이 이해했는지 확인해 보려던 참입니다.

⏱ 29:21설명을 위해 행렬 크기만 11×5로 바꾸고 스레드 블록은 여전히 4×3으로 두었습니다. 처리할 원소는 55개이지만, 작업은 스레드 12개씩의 묶음으로 만들어집니다. numBlocks의 계산에서는 올림을 하므로 실제 생성되는 스레드는 55개보다 많습니다. 각 스레드가 자신의 i와 j를 구하면 일부는 배열 경계 밖의 위치를 가리키게 됩니다. 그래서 코드에 if문을 넣어 스레드가 실제로 유효한 일을 맡고 있는지 검사합니다. 유효한 일이 없는 스레드는 아무것도 하지 않아야 합니다.

⏱ 30:16이 예제는 CUDA가 SPMD 프로그래밍 모델이라는 점을 분명히 보여 줍니다. 스레드들을 생성하면 각 스레드가 프로그램을 통해 자신의 일을 결정합니다. 배열 원소마다 스레드를 하나씩 시작하라는 것이 모델의 정의는 아닙니다. 지정한 모든 블록에 블록당 지정한 수의 스레드를 실행하는 것입니다. 그 안의 코드가 우연히 스레드당 일 하나를 할 수도 있고, 원하는 다른 일을 해도 됩니다. if문을 빼면 어떻게 될까요? 배열 끝을 넘어선 메모리를 읽고 쓰게 됩니다. 결과를 장담할 수 없고 프로그램이 중단될 수도 있습니다.

⏱ 31:03학생이 추상화상으로는 스레드를 만들지만 CPU와 구현 방식이 아주 다른 것이냐고 묻습니다. 맞습니다. 정확히는 CUDA 스레드를 만들며, 하나씩 생성하는 대신 일괄 실행합니다. C++의 스레드 생성 API에 매개변수를 주어 150개를 만들라고 할 수 있다고 상상해 보세요. 여기서는 더 정확히 스레드 블록 150개를 만들고 각 블록은 4×3으로 구성해 달라는 식입니다. 이것이 스레드 일괄 실행입니다. GPU가 이 스레드들을 실행한다는 것 외에는 아직 구현 방법을 설명하지 않았습니다. 다른 학생이 앞서 메모리를 할당하던 슬라이드로 돌아가 달라고 요청합니다.

⏱ 32:03학생은 그 CUDA 작업을 마친 뒤의 모든 CUDA 계산은 deviceA를 대상으로 수행하는 것이냐고 묻습니다. 맞습니다. 그러나 여기의 일반 C++ 코드에서 deviceA[0]을 printf로 출력하려 한다면 세그멘테이션 오류가 나야 합니다. deviceA는 이쪽에서 참조할 수 없는 주소 공간을 가리키는 포인터이기 때문입니다. 저수준 C 코드이므로 이런 잘못된 접근을 시도하는 것 자체는 가능합니다. 따라서 기본 모델에는 CPU 스레드용 주소 공간과 GPU의 CUDA 스레드용 주소 공간이 따로 있습니다. CUDA에는 추가 주소 공간도 있습니다. 스레드 블록이라는 개념을 둔 이유 중 하나가 이와 관련됩니다.

⏱ 33:03각 스레드 블록에는 그 블록에 속한 스레드들만 접근할 수 있는 주소 공간이 있습니다. 개별 스레드에도 자신만 접근하는 로컬 주소 공간과 스택이 있으므로 스레드의 지역 변수는 해당 스레드만 접근합니다. 이 점은 ISPC와 같습니다. 여기에 ISPC의 uniform 변수와 어느 정도 비슷한 블록별 변수, 그리고 어떤 스레드든 load와 store로 접근할 수 있는 디바이스 전역 메모리가 더 있습니다. 스레드 한 묶음이 어떤 일을 함께 하고 다른 묶음은 다른 일을 함께 하며, 협업을 위한 공유 공간을 갖는 구조입니다. 지역성(locality)을 위해 어떤 스레드들을 함께 배치해야 하는지 GPU에 힌트를 주는 셈입니다.

⏱ 34:07CUDA를 연습하기 위해 행렬 덧셈 대신 간단한 1차원 합성곱(convolution)을 해 봅시다. 출력은 해당 위치 주변 입력들의 평균입니다. 영상 처리나 신호 처리의 기본 연산이고, 심층 신경망의 합성곱 층에 해당하는 1차원 형태이기도 합니다. 슬라이드 위쪽에 입력과 출력 위치를 그렸고, 빨간 화살표로 출력 하나를 계산하는 데 필요한 입력 원소들을 표시했습니다. 입력이 출력보다 두 원소 더 크다는 점부터 확인하세요. 경계 조건을 별도로 처리하는 복잡함을 피하고 예제를 단순하게 유지하기 위한 설정입니다.

⏱ 35:01코드에서는 먼저 convolve CUDA 커널을 호출합니다. 배열 크기를 설명하며 처음에는 1,024라고 했다가 약 100만 원소를 예로 들고, 출력 수를 약 100만 개, 즉 1,024²개라고 표현합니다. 정수 계산을 단순하게 유지하려고 N이 블록당 스레드 수로 나누어떨어진다고 가정하고 N/threadsPerBlock개의 블록을 시작합니다. 각 스레드는 자신의 ID를 1차원 배열 인덱스로 해석해 담당 출력을 구하고, 메모리에서 값 세 개를 읽어 곱셈 계산을 수행한 뒤 결과를 씁니다. 이것이 1차원 합성곱의 올바른 CUDA 구현인지 확인해 보세요. 지금까지는 블록당 스레드 개념을 사실상 활용하지 않았으므로, 블록 구문 없이 N개 스레드를 실행한다고만 하고 싶을 정도입니다.

⏱ 36:21각 스레드는 원소 세 개를 읽고, 이웃한 스레드들이 읽는 원소는 서로 겹칩니다. 한 스레드가 어떤 데이터에 load 명령을 내리면 다른 스레드도 같은 데이터를 읽는 것입니다. 캐시 적중 등이 도움이 될 수 있지만, 더 효율적인 프로그램을 작성하는 다른 방법을 보겠습니다. 다음 코드는 훨씬 복잡하지만 블록에 로컬 메모리가 있다는 개념으로 무엇을 할 수 있는지 보여 줍니다. 실행할 일은 같고 스레드 블록 크기는 128입니다. 출력 128개를 계산하려면 입력이 몇 개 필요할까요? 128에 2를 더한 130개입니다.

⏱ 37:18스레드는 원소 130개짜리 공유 배열을 선언합니다. __shared__라는 형식 수식어는 이것이 스레드별 할당이 아니라 스레드 블록별 할당임을 뜻하며, ISPC의 uniform 변수와 비슷합니다. 각 스레드는 인덱스를 계산한 뒤 입력 배열에서 값 하나를 공유 배열로 복사합니다. 첫 스레드는 공유 배열의 인덱스 0에 쓰고 나머지도 차례로 자기 위치를 맡습니다. 이어 블록 내부 threadIdx가 2보다 작은지 검사합니다. 스레드 0과 1은 마지막 두 원소도 가져옵니다. 스레드는 128개인데 입력은 130개가 필요하므로, 모두 하나씩 가져온 뒤 그중 두 스레드가 일을 조금 더 하는 것입니다. 코드가 보기 좋지는 않지만 역할을 명시적으로 나눈 것입니다.

⏱ 38:51잠깐 한 줄을 건너뛰어 보면 각 스레드는 원래 하던 합성곱을 계산하되, 주 전역 메모리 대신 이 공유 변수를 읽습니다. 남은 한 줄인 __syncthreads()는 스레드 블록 안에서 사용하는 CUDA 장벽(barrier)입니다. SPMD로 작성된 스레드들은 동시에 진행될 수 있습니다. 모든 스레드가 협력하여 공유 배열에 데이터를 적재하고 장벽에 도달하면, 블록의 모든 스레드가 자기 몫의 일을 끝낼 때까지 기다립니다. 필요한 데이터가 전부 들어왔음을 확인한 뒤에는 각자 출력 원소를 계산하며 다시 독립적으로 병렬 실행합니다.

⏱ 39:45왜 계산을 이렇게 세심하게 조율할까요? 구현 관점에서 이 공유 변수는 고성능 L1 캐시처럼 생각할 수 있는 저장 공간을 사용합니다. CUDA는 같은 블록의 스레드들을 대체로 같은 코어에 배치하므로, 재사용되는 데이터에 이 빠른 로컬 메모리로 접근할 수 있습니다. 여기서는 스레드 하나가 같은 데이터를 두 번 읽는 것은 아닙니다. 한 스레드가 쓰는 데이터를 이웃 스레드도 쓴다는 점에서 재사용이 생깁니다. 그래서 협력하여 데이터를 한 번만 가져오고, 그 뒤에는 아주 빠른 공유 메모리 또는 로컬 메모리에서 작업하는 것입니다.

⏱ 40:42학생이 왜 __syncthreads()가 필요한지, 작은 격자의 데이터를 적재한 다음 계산하면 안 되는지 묻습니다. 질문 일부는 들리지 않습니다. 함께 생각해 봅시다. 이 호출을 없애면 왜 프로그램이 틀릴 수 있을까요? 다른 학생이 어떤 스레드가 접근할 데이터가 아직 공유 공간에 적재되지 않았을 수 있다고 답합니다. 맞습니다. GPU가 코드를 어떤 순서로 실행할지 보장되지 않습니다. 한 스레드가 자기 값을 적재하자마자 계산을 시작했는데 다른 스레드들은 아직 값을 가져오지 않았을 수 있습니다. 장벽을 통과한 뒤에는 공유 로컬 메모리가 올바르게 초기화되었다는 보장이 생깁니다.

⏱ 41:32다른 학생이 별도의 스레드용 공간이 왜 필요한지, 큰 블록 메모리 하나를 쓰면 안 되는지 묻습니다. result 같은 지역 변수를 말하는 것이군요. 프로그래밍 모델에는 당연히 스레드별 지역 변수가 필요합니다. 기본적으로 모든 변수는 스레드마다 사본이 있고, 추가로 블록당 한 번 할당하는 변수도 만들 수 있습니다. 실제로 어느 메모리에 배치할지는 컴파일러가 결정합니다. 모든 스레드를 위한 큰 메모리 블록 안에 배치할 수도 있습니다. 다만 개념적으로 자기 스레드의 실행에만 속하는 비공개 변수와 다른 스레드와 나누는 변수가 둘 다 필요하다는 것은 자연스럽습니다. 아직 물리적 배치에 지나치게 신경 쓸 필요는 없습니다.

⏱ 42:30학생이 이 스레드들이 SPMD 방식으로 실행되느냐고 묻습니다. CUDA 코드는 SPMD가 맞습니다. 이어지는 질문에서는 SPMD와 SIMD를 혼동하고 있습니다. SPMD는 구현과 독립적인 프로그래밍 모델의 문제입니다. ISPC와 CUDA 모두 프로그램 하나를 작성하면 시스템이 서로 다른 스레드 ID를 부여해 여러 번 실행합니다. 그 스레드나 인스턴스들을 효율적으로 실행하는 방법은 구현이 결정합니다. ISPC는 SIMD 명령어를 사용합니다. GPU의 구현은 아직 설명하지 않았습니다. 물론 SIMD가 들어가지만 방식에는 조금 차이가 있습니다.

⏱ 43:24학생이 이것을 프로그래머가 결정하느냐고 묻습니다. 전적으로 프로그래머가 제어합니다. C++에서 실행할 스레드 수를 스스로 정하듯 여기서도 스레드 수를 정하고, 블록으로 어떻게 묶을지도 시스템에 알려 줍니다. 이어 블록 크기가 언제나 코어 하나가 실행할 수 있는 스레드 수와 같아야 하는지 묻습니다. 코어 하나가 실행할 수 있는 수보다 더 많은 스레드를 블록 하나에 넣을 수는 없습니다. 그러나 더 적게 넣는 것은 가능합니다. 코어 하나에 블록 여러 개를 함께 배치할 수 있기 때문입니다.

⏱ 44:03학생이 이 예제처럼 인접한 데이터를 읽는 대신 행렬의 열을 읽을 때도 같은 블록 메모리 방식을 쓰는 것이 좋은지, 블록 안에서 연속 주소를 읽는지가 중요한지 묻습니다. 성능에는 분명히 중요합니다. 여기서는 모든 스레드가 연속 주소에 접근하므로 내부 구현에 유리하리라고 예상할 수 있습니다. 임의의 주소에서 값을 가져와 공유 메모리에 넣어도 프로그램은 동작하지만, 캐시 지역성 등 메모리 시스템의 특성 때문에 성능이 낮아질 수 있습니다. cudaMalloc이 데이터를 배치하는 방식도 제어할 수 있느냐는 질문에는, C++의 malloc처럼 연속된 주소 공간을 할당할 뿐이라고 답합니다. 그 안에 행렬 원소들을 어떻게 배치할지는 작성자에게 달렸습니다. 여기서 연속적인 행 우선(row-major) 배치를 가정한 것도 제 선택입니다.

⏱ 45:25ISPC에는 인스턴스 사이의 통신을 위한 레인 간 연산이 있었습니다. CUDA에는 여러 동기화 수단이 있습니다. 스레드 블록별 장벽, 원자적 연산(atomic operation), 그리고 CUDA 스레드를 시작하고 끝날 때까지 기다리는 것과 같은 호스트와 디바이스 사이의 동기화가 있습니다. 과제 3에서 이런 것들을 만나게 됩니다. CUDA는 스레드 ID, 블록 ID 등에 따라 스레드가 하는 일을 하나의 프로그램으로 정의하는 SPMD 모델입니다. 필요하면 같은 블록의 스레드끼리 장벽으로 동기화할 수 있고, 실행할 때는 많은 스레드를 한꺼번에 시작합니다. 이제 이 모델의 구현을 생각해 봅시다.

⏱ 46:26이 프로그램은 앞의 convolve와 완전히 같고 블록 크기는 128입니다. 호스트의 C++ 코드에서는 약 100만 스레드를 만들며, 이를 128로 나누면 약 8K개의 블록입니다. NVIDIA는 이를 어떻게 구현할까요? 의미상으로는 CUDA 스레드 100만 개를 실행해야 하지만, 실제로 하드웨어 스레드 문맥 100만 개를 찾아 배정할까요? 아마 그런 수의 실행 문맥 자체가 없을 것입니다. 이 대량 실행은 무엇과 비슷한가요? 학생이 태스크라고 답합니다. 맞습니다. 많은 일, 즉 이 모든 스레드 블록을 줄 테니 처리하라는 것입니다. ISPC 태스크가 프로그램 인스턴스들로 나뉘듯 CUDA 스레드 블록도 CUDA 스레드들로 나뉩니다.

⏱ 47:40사용자가 코드를 고치지 않아도 코어 수나 SIMD 폭과 관계없이 실행할 수 있어야 합니다. 코어가 많은 큰 GPU, 코어가 적은 중급 GPU, 요즘 얼마인지 모르지만 10만 달러쯤 한다는 그레이스 호퍼(Grace Hopper) 칩에서도 같은 코드가 동작하기를 바랍니다. 호스트 CPU 코드와 디바이스 코드를 다시 보세요. 디바이스 코드를 컴파일하면 명령어 흐름뿐 아니라 실행에 필요한 조건을 담은 메타데이터도 생성됩니다. 프로그래머가 블록당 스레드 128개를 요구했으므로, 컴파일된 프로그램은 그만큼의 스레드 자원이 필요하다는 사실을 알 수 있습니다.

⏱ 48:29블록당 스레드가 128개이면 공유 배열은 원소 130개가 필요합니다. 따라서 float 130개를 저장해야 합니다. 설명 중 처음에는 512바이트라고 말하지만, 이어 필요한 저장 공간을 520바이트로 바로잡습니다. 결국 이 블록을 실행하려면 스레드 128개분의 자원과 공유 저장 공간 520바이트가 필요합니다. 세부 구조는 잠시 잊고 코어가 네 개인 작은 GPU를 떠올려 보세요. 위에서 약 8,000개의 스레드 블록을 만들었습니다. 스레드 풀을 다뤄 본 사람이라면 모든 태스크, 즉 블록들을 여러 작업자에게 차례로 나누어 주는 방식임을 이해할 수 있을 것입니다. CUDA에서는 이 스레드 블록들 사이에 의존성을 설정하는 것도 가능합니다.

⏱ 49:29과제 2에서는 어느 작업자 스레드에 어떤 일을 줄지 결정하는 소프트웨어를 여러분이 작성합니다. GPU에서는 그 기능이 실리콘의 하드웨어에 들어 있습니다. 과제에서 구현하는 ‘이 태스크를 이 횟수만큼 실행하라’는 API가 이제 하드웨어 인터페이스인 셈입니다. 프로그램과 프로그램 카운터를 주고 특정 스레드에서 실행을 시작하라는 방식이 아닙니다. 명령어 흐름, 즉 프로그램 바이너리를 주고 그 프로그램의 모든 스레드 블록을 실행하라고 요청합니다. GPU 상단의 하드웨어는 블록 8,000개와 코어 네 개가 있음을 보고 작업을 시작합니다. 독립적인 일을 선언하고 스케줄러가 작업자에게 배정하게 하는 설계 패턴의 또 다른 사례입니다.

⏱ 50:26이제 작업자들을 살펴보겠습니다. 볼타(Volta) V100의 구조 그림을 사용합니다. 이번 주 후반 제 개인 마감이 끝나고 여러분이 CUDA 과제를 시작할 때쯤, 흥미 삼아 현대 그레이스 호퍼 칩이 어떻게 생겼는지도 슬라이드 몇 장으로 만들 수 있을 것입니다. 기본적인 모습은 꽤 비슷하고 이런 아이디어는 크게 변하지 않았습니다. 작은 GPU 프로세서, 즉 GPU 코어를 생각해 보세요. 위에는 익숙한 명령어 인출·해독(fetch and decode) 장치가 있습니다. 무지개색 표시는 일단 무시하세요. 옆의 주황색 상자들은 SIMD 방식으로 동작해야 하는 ALU 16개, 즉 실행 장치 16개입니다.

⏱ 51:14여기에는 실행 문맥을 저장하는 공간이 있습니다. 이렇게 그린 이유는 스레드 0에 r0, r1, r2, r3가 있고, 스레드 1과 2, 31에도 각각 대응하는 저장 공간이 있으며 그 뒤로도 계속 이어진다는 점을 보여 주기 위해서입니다. 이 그림에는 스레드 128개분의 실행 문맥이 있습니다. 멀티스레딩을 많이 사용하는 칩입니다. 지금까지 그린 구조에는 슈퍼스칼라는 없고, 스레드 문맥 128개와 폭 16인 SIMD 실행 장치가 있습니다. 다른 장치들은 부동소수점 SIMD, 정수 SIMD, 폭 8의 삼각함수 같은 특수 수학 연산, load와 store 등 여러 종류의 SIMD 연산을 지원합니다. NVIDIA는 이 구조를 SM 서브코어라고 부릅니다.

⏱ 52:35스레드 하나에 딸린 이 레지스터들은 스칼라 레지스터입니다. CPU의 SIMD를 다룰 때는 벡터 레지스터라고 했지만, 여기의 레지스터는 벡터가 아닙니다. NVIDIA 스레드의 실행 문맥은 스칼라 레지스터들의 집합이며, 스레드 2에도 별도의 스칼라 레지스터 집합이 있습니다. 여기서 CPU와 GPU에서 SIMD가 나타나는 방식의 큰 차이가 나옵니다. CPU 스레드가 SIMD를 실행할 때는 컴파일러가 SIMD 명령어를 생성합니다. 반면 GPU는 하드웨어에 대한 인터페이스 자체가 SPMD이므로, 같은 프로그램을 N개 실행한다는 사실을 하드웨어가 압니다.

⏱ 53:30각 스레드는 자기 프로그램 카운터(PC)를 레지스터 중 하나로 가집니다. 연속된 스레드 32개의 프로그램 카운터가 같으면 그 명령어를 SIMD 방식으로 함께 실행할 수 있습니다. 기본 아이디어는 같지만 구현상의 차이가 있습니다. CPU에서는 컴파일러가 SIMD 명령어를 만들고, GPU에서는 같은 프로그램을 실행하는 CUDA 스레드들이 같은 위치에 도달했을 때 하드웨어가 SIMD 장치로 한꺼번에 실행합니다. 이를 암묵적 SIMD(implicit SIMD)라고 부릅니다. 결국 같은 병렬 실행 아이디어의 다른 구현입니다.

⏱ 54:18아주 최근까지, 특히 이런 오래된 GPU에서는 ID가 연속인 스레드들이 같은 연산을 할 때만 SIMD로 묶을 수 있었습니다. NVIDIA가 여기 스레드 하나와 저기 스레드 하나를 임의로 가져와 SIMD로 실행하지는 않았습니다. 따라서 실행 문맥 32개를 묶은 것은 벡터 레지스터가 있는 CPU 스레드 하나처럼 생각할 수 있습니다. NVIDIA에서는 이를 워프(warp)라고 부릅니다. 워프는 CUDA 스레드 32개분의 실행 문맥이며, 그 스레드들이 같은 일을 할 때 SIMD와 같은 실행을 받습니다. 같은 프로그램 지점에 있는지는 하드웨어가 기본적으로 PC들을 비교해 확인하고, 공통 명령어를 발행합니다.

⏱ 55:41학생이 워프는 32개 스레드라면서 SIMD ALU는 16개뿐이라는 점을 짚습니다. 실제로는 연속된 두 클럭에서 16개씩, 다시 16개씩 실행합니다. 그래서 워프의 해당 연산은 두 클럭이 걸립니다. 그렇다면 명령어 인출·해독 장치가 한 클럭씩 놀게 될까요? 그렇지 않습니다. 남는 클럭에는 다른 실행 장치로 명령어를 보내기 시작합니다. 이렇게 하면 인출·해독 장치를 덜 두고도 됩니다. 폭 32인 연산을 발행하면 그 실행 장치를 다시 채우기까지 두 클럭이 있으므로, 지금 부동소수점 연산을 보내고 다음 클럭에 정수 연산을 보낸 뒤 다시 부동소수점으로 돌아오는 식입니다.

⏱ 56:35시간축에는 클럭과 명령어 0, 1, 2가 표시됩니다. 두 사이클 걸리는 워프 명령어를 보내고, 다음에는 모든 스레드의 i++ 같은 정수 명령어를 보내며, 두 사이클 뒤 FP32 장치에 다시 명령어를 보낼 수 있습니다. 수업의 핵심은 아닌 저수준 구현 세부 사항이지만 흥미로울 것 같아 소개합니다. 학생이 스케줄러가 명령어를 해독하고 실행 장치 종류와 수를 파악해 적절한 장치로 보낼 버퍼를 관리할 만큼 똑똑하냐고 묻습니다. 그렇습니다. 어떤 의미에서는 일반적인 비순차 실행(out-of-order) 파이프라인 프로세서의 기능입니다. 다만 아주 동적인 것은 아닙니다. 부동소수점 명령어를 발행하면 다음 것은 두 사이클 뒤에 보낼 수 있다는 정도로 단순하며, CPU 수준으로 정교하지 않습니다.

⏱ 57:38서로 다른 장치 묶음에는 부동소수점 ALU, 8비트 정수 장치, 32비트 정수 장치, load/store 장치, 초월함수 연산 장치 등이 있습니다. 인출·해독 장치가 CUDA 프로그램을 따라가며 이 32개 스레드 묶음에 필요한 부동소수점 또는 정수 연산을 보내는 것입니다. 세부 사항이 아주 중요한 것은 아닙니다. 또 다른 질문에 답하자면 모든 스레드에는 실제로 자기 PC가 있지만, 워프 안의 PC가 같을 때만 SIMD 실행을 얻습니다. 그래서 앞서 128개 스레드를 그린 그림에서는 사실상 PC가 네 개뿐인 것처럼 보일 수 있습니다.

⏱ 58:26그러나 이 추상화는 NVIDIA에 몇 가지 선택지를 남겨 줍니다. 원한다면 스레드를 재구성하여 SIMD의 실행 일치성을 높일 수 있습니다. 제 생각에는 아주 최근에야 매우 제한적인 조건에서 그런 일을 하기 시작했습니다. 어떤 의미에서는 99%의 경우 폭 32의 SIMD 기계가 서로 다른 PC 네 개를 가진 것처럼 동작합니다. 워프를 벡터 명령어를 쓰는 보통 스레드처럼 생각하면 유용하지만, 그것이 문자 그대로의 구현은 아닙니다. 실제로는 CUDA 스레드마다 서로 다른 PC가 있고, 칩이 벡터 실행을 하기 전에 그 값들이 같은지 동적으로 확인합니다.

⏱ 59:19그림을 바꾸는 중 실행 문맥을 60워프분으로 확장했다고 언급하지만, 다시 단순한 그림으로 돌아와 여기에는 16워프가 있다고 설명합니다. 실행 문맥은 16×32개이고 PC도 그 수만큼 있으며, 이 서브코어에서는 한 번에 그중 한 워프의 명령어 흐름을 선택해 실행합니다. 학생이 워프도 스레드의 집합인지, 스레드 블록과는 무엇이 다른지 묻습니다. 워프는 하드웨어 구현의 세부 사항이고 블록은 프로그래밍 모델의 일부입니다. 프로그래머는 CUDA 스레드 128개를 요청했을 뿐이고, NVIDIA는 그것을 실행하기 위해 실행 문맥 네 워프분을 사용합니다.

⏱ 1:00:25전체를 합쳐 보겠습니다. 이 내용은 시험에 내지 않고, 세부 구조를 살펴보는 것이 재미있어서 설명하는 것입니다. 실제 코어에서는 이 서브코어가 네 번 복제됩니다. 실행 문맥에 그런 번호를 붙인 이유도 이것입니다. 큰 코어 전체에는 64워프, 즉 64×32=2,048개로 약 2,000개의 CUDA 스레드 실행 문맥이 있습니다. 벡터 ALU 묶음 네 개와 인출·해독 장치 네 개가 있고, 전체에 공유 메모리가 연결됩니다. 폭 4의 슈퍼스칼라처럼 생각하면 매 클럭 각 인출·해독 장치가 실행할 워프를 하나씩 찾습니다. 실제로는 각 장치가 자신에게 배정된 워프 집합 안에서만 선택합니다.

⏱ 1:01:33네 개의 별도 코어처럼 보이지만, 모든 실행 문맥이 같은 공유 메모리 저장 공간을 나누므로 완전히 별개는 아닙니다. 앞의 질문과 연결해 약 2,000개 CUDA 스레드와 공유 저장 공간 128을 요구하는 블록을 상상해 봅시다. 여기서 128의 단위는 원문에 명시되지 않았습니다. 그런 블록이 이 구조 전체를 차지하는 식입니다. 이 코어를 SM이라고 합니다. 약 2,000개 CUDA 스레드 사이에서 멀티스레딩하며, 앞서 말한 구획 제약 아래 사용 가능한 64워프 중 최대 네 워프를 골라 매 클럭 진행시킬 수 있습니다. 지연 시간을 많이 숨길 수 있습니다. 여전히 새로운 개념이 아니라 멀티스레딩, SIMD, 동시 멀티스레딩, 명령어 인출 장치들에서 나오는 슈퍼스칼라 발행의 조합입니다.

⏱ 1:02:26학생이 한 번에 16개를 실행할 수 있다는 뜻인지 확인합니다. 아닙니다. 현재 슬라이드만 보면 각 인출·해독 장치가 워프 하나를 고르므로 한 클럭에 네 워프를, 장치마다 하나씩 발행합니다. 동시에 진행시킬 수 있는 워프는 네 개지만 선택할 후보는 64개입니다. 인텔 프로세서가 64개의 하이퍼스레드 중에서 골라 네 스레드의 벡터 명령어를 동시에 실행하는 상황과 비슷합니다. 다른 학생은 워프를 골라도 같은 PC를 가진 일부 스레드만 실행될 수 있으니 워프 전체가 실행되는 것은 아니라고 짚습니다. 맞습니다. 워프를 선택하고 PC가 같은 스레드들을 실행하며 나머지는 익숙한 SIMD 마스크 방식으로 가립니다.

⏱ 1:03:27학생이 계속 나오는 PC가 무엇인지 묻습니다. 프로그램 카운터(program counter)이며, 다음에 실행할 명령어를 가리키는 레지스터입니다. 하드웨어는 워프의 스레드들이 같은 프로그램 지점에 있고 실행 가능한지를 보고 SIMD ALU로 보냅니다. 이 세부 사항에 너무 오래 머무르지 않고 넘어가겠습니다. 다른 학생이 명시적으로 고정한 SIMD에 비해 이런 방식의 장점이 무엇인지 묻습니다. 동작은 익숙한 SIMD와 같습니다. 워프의 모든 PC가 같지 않으면 같은 것들끼리 먼저 실행하고, 다른 PC의 스레드들을 위한 명령어를 따로 실행합니다. 과제 2의 마스킹과 비슷하지만 무엇을 함께 실행할지 컴파일 시점에 확정하는 대신 하드웨어가 결정합니다.

⏱ 1:04:30이 방식으로 NVIDIA는 미래 칩의 SIMD 폭을 예를 들어 64로 바꿀 가능성을 열어 둡니다. 인텔이 SIMD 폭을 바꾸면 바이너리에 ‘폭 8인 AVX 명령어’처럼 구체적인 내용이 들어 있으므로 코드를 다시 컴파일해야 합니다. NVIDIA는 프로그램이 스칼라 명령어들로 이루어져 있어 폭을 바꾸더라도 기존 코드가 계속 실행되기를 원합니다. 스레드가 서로 다른 경로로 갈 때 내부 연산 순서를 바꾸어 실행 일치성을 높일 여지도 원하고, 실제로 일부 경우에는 조금 그렇게 합니다. 그래도 대부분의 코드는 평범한 SIMD로 생각하면 현대 GPU의 동작을 잘 이해할 수 있습니다.

⏱ 1:05:15학생이 블록 크기를 사용자가 정할 수 있다고 했는데 NVIDIA 아키텍처가 바뀌면 어떻게 되느냐고 묻습니다. 맞습니다. 프로그램에 어떤 요구를 넣으면 하드웨어가 그것을 지원해야 한다고 선언하는 셈입니다. 이 특정 프로세서 그림을 놓고 가정해 봅시다. 여기에는 CUDA 스레드 2,048개가 들어가고 같은 공유 메모리에 접근할 수 있으니, 스레드 2,048개짜리 블록을 만들 수 있다고 해 보는 것입니다. 이것은 해당 그림을 이용한 가정입니다. 이전 GPU들에는 실행 문맥 수의 최댓값이 있었습니다.

⏱ 1:05:46이전 GPU는 SM 코어당 실행 문맥이 최대 256개였습니다. 블록당 약 2,000개 스레드를 요구하는 프로그램은 그런 GPU에서 실행할 수 없고 지원 한도가 256개라는 이유로 거절됩니다. 학생이 이전 아키텍처에 맞추어 쓴 프로그램은 새 구조에 자동으로 확장되지 않는 것이냐고 묻습니다. 맞습니다. 블록당 256개로 작성한 뒤 약 2,000개를 지원하는 새 GPU가 나오면 프로그램을 바꿀 수도 있고 그대로 둘 수도 있습니다. 그대로 두면 NVIDIA가 코어당 블록 네 개를 실행할 수도 있습니다. 수동으로 스케줄링할 때마다 하드웨어 능력에 대한 가정을 하게 됩니다. NVIDIA는 묶음 크기를 바꾸거나 연산 순서를 조정할 수 있으므로 SIMD 스케줄링을 자신에게 맡기기를 원하고, 때로는 더 잘할 수도 있습니다.

⏱ 1:06:54이 절을 마무리하며 블록당 스레드 128개인 코드로 돌아갑니다. 여기서는 저장 공간을 다시 512바이트라고 설명합니다. CUDA는 블록을 코어에 올릴 때 그만큼의 공유 메모리와 실행 문맥 128개, 즉 네 워프를 할당합니다. 워프들을 여러 슬라이스에 배치하여 동시에 진행할 수 있도록 합니다. 그림의 전체 구조가 NVIDIA의 스트리밍 멀티프로세서(streaming multiprocessor), 즉 SM이고, V100에는 이것이 80개 반복됩니다. 일반 클럭을 약 1.2GHz로 보고, SM 80개와 SM당 네 열, 열마다 폭 16의 부동소수점 벡터 ALU를 함께 계산하여 부동소수점 성능을 12.7테라플롭스라고 설명합니다.

⏱ 1:08:1280개 SM에 SM당 64워프, 워프당 32스레드를 곱하면 칩에 동시에 살아 있는 CUDA 스레드가 약 163,000개입니다. 여기서 비교하는 인텔 프로세서는 하이퍼스레드 두 개를 지원하는 반면 GPU는 약 163,000개 CUDA 스레드를 지원합니다. 모두 같은 순간에 명령어를 실행하는 것은 아니지만, 살아 있는 실행 문맥이 동시에 칩에 존재합니다. 학생이 각 스레드의 메모리가 공유 메모리에 배정된다는 뜻인지 묻습니다. 공유 메모리에 들어가는 것은 공유 변수입니다. 프로그램의 shared float support 선언은 스레드 블록당 한 번 이루어지는 할당입니다.

⏱ 1:08:53공유 메모리는 프로그래머가 실제 스케줄링 세부 조건을 요구하기 시작하는 지점입니다. CUDA는 절충합니다. 블록 수만 알려 주면 실행 방법은 알아서 정하지만, 메모리는 너무 중요하므로 여기서처럼 저장 공간 512바이트가 필요하다고 명시할 수 있게 합니다. NVIDIA는 블록을 실행할 때 같은 코어의 워프 실행 문맥을 선택하고, 그 문맥들이 모두 읽고 쓸 수 있는 저장 공간을 할당합니다. 스레드들이 같은 코어에 있으므로 장벽도 빠릅니다. 블록의 일부 스레드를 컴퓨터의 다른 위치에 두면 장벽을 위해 여러 위치 사이의 통신이 필요해집니다. 따라서 스레드 블록은 프로그램에서 동기화와 데이터 지역성의 단위입니다.

⏱ 1:10:03여기서는 단순화한 CUDA를 설명하고 있습니다. 현대 CUDA 문서에는 공유 메모리에서의 malloc을 비롯한 추가 할당 방식도 나오지만, 그런 내용을 넣으면 더 복잡해지므로 먼저 단순한 형태를 다룹니다. 이제 블록마다 512바이트가 필요한 이 프로그램을 실행한다고 합시다. 방금 본 복잡한 GPU 대신 아주 단순한 GPU를 가정합니다. 클럭당 명령어 하나를 실행하고 실제 폭 32의 SIMD를 가지며, SM 코어마다 CUDA 스레드 384개, 즉 12워프분의 실행 문맥이 있습니다. 어떻게 실행될까요?

⏱ 1:10:48약 8,000개 블록을 실행하라고 요청합니다. 큐나 스레드 풀을 구현 중인 여러분에게 익숙한 패턴입니다. convolve 태스크와 인수, 블록 수를 주는 것이며, 작업 기술자의 예를 말할 때는 블록 1,000개라는 수도 언급합니다. GPU 작업 스케줄러는 첫 블록에 실행 문맥 128개와 공유 메모리 512바이트를 배정하고 다음 포인터를 1로 옮깁니다. 다음에는 무엇을 할까요? 더 많은 일을, 가능하면 다른 코어에 배치합니다. 두 번째 블록에도 실행 문맥 128개와 512바이트를 주고 다음 포인터는 2가 됩니다. 계속해서 한 프로세서에 여러 블록을 배치할 수 있으며, 블록마다 문맥 128개와 512바이트를 추가로 배정하여 네 번째 블록까지 채웁니다.

⏱ 1:12:02이제 다섯 번째 블록은 자원이 부족해 당장 배치할 수 없습니다. 이 예에서는 실행 문맥은 남아 있지만 공유 메모리가 부족하며, 공유 메모리가 1.5K뿐이라고 설명합니다. NVIDIA는 들어갈 수 없음을 아는 블록을 배치하지 않고 기다립니다. 어느 블록이 끝나 자원이 풀리면 다음 블록인 인덱스 4번을 배치합니다. SPMD 프로그램이고 모든 블록의 자원 요구량이 같으므로, 블록 하나가 끝날 때 새 블록 하나에 필요한 자원이 정확히 확보됩니다. 요구량이 균일하여 복잡한 할당 문제를 풀 필요가 없습니다. 자원 요구량을 미리 아는 것이 편리한 이유이며, 이 과정을 모든 블록을 처리할 때까지 계속합니다.

⏱ 1:13:04남은 4~5분 동안 이해를 확인해 봅시다. 다음 슬라이드 세 장을 이해하면 내부 구현을 머릿속으로 단계별 실행해 볼 수 있을 것입니다. 실행 문맥이 128개, 즉 네 워프밖에 없는 작은 코어와 블록당 CUDA 스레드 256개를 요구하는 프로그램을 상상하세요. NVIDIA는 이 실행을 거절합니다. 왜 처음 128개를 끝까지 실행한 뒤 다음 128개를 실행하면 안 될까요? 학생이 프로그래머가 데이터 지역성을 지정했기 때문이라고 답합니다. 프로그래밍 모델의 규칙으로는 맞지만, 왜 규칙을 그렇게 정했을까요? 공유 메모리는 그대로 제공하면서 임의 크기의 블록을 앞 절반, 뒤 절반으로 실행해 주면 안 될까요?

⏱ 1:14:35다른 답은 프로그래머가 어떤 실행이 효율적인지 알고 있다는 것입니다. 시스템이 프로그래머의 요청을 존중하는 것은 좋지만, 구현자가 정확성을 보존하면서 더 나은 방법을 택하는 경우도 많습니다. 여기서는 프로그램 안의 장벽을 보세요. 모든 스레드가 도착해야 그다음으로 진행할 수 있습니다. 첫 128개를 실행하면 장벽에서 멈춘 채 실행 문맥을 계속 차지하고 그것을 내놓지 않습니다. 나머지 128개는 실행 기회를 얻지 못합니다. 그러면 아무 일도 더 진행되지 않고 프로그램은 교착 상태(deadlock)에 빠집니다.

⏱ 1:15:21이를 피하려면 더 복잡한 구현이 필요합니다. 상주한 스레드를 선점하여 프로세서에서 빼고 다른 스레드를 올려야 하며, 성능이 매우 나빠질 수 있습니다. NVIDIA는 고성능 프로그래밍 시스템을 제공하므로 모든 요청을 어떻게든 정확하게 실행시키려고 성능을 해치는 복잡한 장치를 넣고 싶지 않습니다. 여기서 다루는 표준 CUDA 모델에서는 선점이나 마법 같은 개입이 없다고 생각하면 됩니다. 그래서 이 프로그램을 거절합니다. 같은 블록의 스레드는 단지 SPMD 스레드인 데 그치지 않고, 장벽이나 원자적 연산으로 협력할 수 있도록 같은 SM에서 동시에 살아 있는 상태여야 합니다. 따라서 너무 큰 블록을 요구한 이 예는 잘못된 것입니다.

⏱ 1:16:22예를 두 개 더 보겠습니다. 첫째는 히스토그램을 만드는 CUDA 프로그램입니다. 여러 블록에 속한 스레드들이 각자 값을 계산하고 해당 빈(bin)의 값을 증가시킵니다. CUDA에는 atomicAdd 같은 원자적 연산이 있습니다. 우리 모두가 CUDA 스레드라고 상상해 보세요. 각자 값을 가져와 함수를 계산하고 ‘42번 빈을 증가시켜야 한다’고 결정하면, 그 빈에 가서 안전한 원자적 갱신을 수행합니다. 서로 다른 블록의 스레드들이 같은 메모리에 접근하는데 유효한 프로그램일까요? atomicAdd가 전역 메모리에 대해서도 원자성을 제공한다고 가정하면 유효합니다.

⏱ 1:17:22이 프로그램은 CUDA가 블록들을 어떻게 실행할지에 대해 아무 가정도 하지 않습니다. CUDA는 원하는 때에 각 블록을 실행할 수 있고 원자적 갱신에 경합이 생길 수는 있지만 괜찮습니다. 필요한 만큼 블록을 만들고 어떤 순서로 실행해도 허용하는 것입니다. 이를 두 블록을 실행하면서 하나는 메모리의 변수를 갱신하고 다른 하나는 그 변수가 갱신될 때까지 기다리는 코드와 비교해 보세요. 한 블록이 자신이 일을 했다고 알리고 다른 블록이 그 알림을 기다립니다. 이제 GPU가 한 번에 블록 하나만 실행할 수 있다고 상상해 봅시다.

⏱ 1:18:10만약 기다리는 쪽인 블록 1을 먼저 실행하면, 블록 0은 실행 기회를 얻지 못하고 블록 1은 영원히 기다립니다. 이 차이를 이해해야 합니다. 블록 사이의 통신 자체는 가능하지만 어떤 순서로 실행되는지 가정할 수는 없습니다. 블록들은 상호작용할 수 있어도 실행 순서를 전제해서는 안 됩니다. 반면 같은 블록 안에서는 스레드들이 동시에 살아 있다는 전제 아래 장벽을 사용할 수 있습니다. 이제 여러분이 다음 일정으로 갈 수 있도록 여기서 마치겠습니다. 질문이 있으면 자유롭게 해도 됩니다.