간단한 CUDA 벡터 덧셈 예제를 따라가며, 컴파일부터 드라이버, 푸시버퍼, QMD, 워프 스케줄링, 메모리 계층, 그리고 결과가 CPU로 돌아오기까지 GPU 커널 실행의 전체 경로를 살펴봅니다.
간단한 CUDA 프로그램이 있다. 두 벡터를 더한다.
__global__ void vadd(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];
}
int main() {
int n = 1 << 20; // a million floats (1,048,576)
size_t bytes = n * sizeof(float);
float *a = (float*)malloc(bytes), *b = (float*)malloc(bytes),
*c = (float*)malloc(bytes);
for (int i = 0; i < n; i++) a[i] = b[i] = 1.0f;
float *da, *db, *dc;
cudaMalloc(&da, bytes);
cudaMalloc(&db, bytes);
cudaMalloc(&dc, bytes);
cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice);
vadd<<<4096, 256>>>(da, db, dc, n); // 4096 * 256 = n threads, one per float
cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost);
printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]);
}
RTX 4090용으로 컴파일하고 실행하면, 라는 사실을 백만 번 올바르게 계산한다. 전부 다 확인해 보지는 않았지만..
$ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd
c[0]=2.000000 c[n-1]=2.000000
이 사실을 알려주기까지 수천만 개의 CPU 명령, 몇 개의 디바이스 파일, 900번의 ioctl, 그리고 메모리 매핑된 도어벨 레지스터 하나가 동원된다. 이 글에서는 이 커널 하나를 코드에서 워프까지 내려갔다가, 다시 답으로 올라오는 과정을 따라간다. 덧붙이자면, 이 글은 에이전트가 만들어낸 ‘가독성 전환’의 한 사례이기도 하다. 호기심과 (기계로 강화된) 끈기만 있으면 컴퓨터에 관해 알아낼 수 없는 것은 정말 거의 없다. 가독성이 AI가 우리가 무엇을 알도록 도울 수 있는지에 대해 갖는 함의에 대한 흥미로운 논의는 여기에서 볼 수 있다..
nvcc로 프로그램 컴파일하기§이 CUDA 프로그램을 디바이스가 실제로 읽을 수 있는 무언가로 바꾸는 방법부터 시작해야 한다. 그러려면 컴파일러가 필요하다. 정확히는, 많은 컴파일러가 필요하다.
nvcc는 여러 다른 컴파일러를 실행하고 그 출력을 결합하는 드라이버 프로그램이다. --keep를 넘기면 전체 파이프라인을 디스크에 남겨 두어서 직접 읽어볼 수 있다.
$ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls
...
vadd.ptx # device code as PTX (from cicc)
vadd.sm_89.cubin # device code as SASS (from ptxas)
vadd.fatbin # cubin + PTX, bundled (from fatbinary)
vadd.cudafe1.stub.c # host launch stub + kernel registration
vadd.o # final host object, fatbin embedded
...
호스트 코드는 호스트 컴파일러로 간다. 디바이스 코드(vadd)는 더 많은 단계를 거친다. LLVM 기반 컴파일러인 cicc가 이것을 PTX로 바꾸고, 이어서 ptxas가 PTX를 SASS로 바꾼다.
PTX는 가상 ISA다. 타입이 붙은 레지스터가 무한히 많고, 하드웨어가 실제로 몇 개를 갖고 있는지에 대한 개념이 없다. 아래는 PTX에서의 vadd 본문(일부 생략)이다.
$ cat vadd.ptx
...
mad.lo.s32 %r1, %r3, %r4, %r5; // set register r1 to ctaid*ntid + tid
setp.ge.s32 %p1, %r1, %r2; // set predicate p1 if i >= n
@%p1 bra $L__BB0_2; // if out of bounds, skip to exit
cvta.to.global.u64 %rd4, %rd1; // convert generic pointer %rd1 to a global address, store in %rd4
mul.wide.s32 %rd5, %r1, 4; // multiply r1 by 4, store the result in %rd5
add.s64 %rd6, %rd4, %rd5; // add %rd4, %rd5, result in %rd6
ld.global.f32 %f2, [%rd6]; // load a[i] into %f2
...
add.f32 %f3, %f2, %f1; // add %f1 and %f2, result in %f3
st.global.f32 [%rd10], %f3; // store c[i] = ... in global memory
가상 레지스터는 %rd1–%rd10, %f1–%f3처럼 보인다. 접두사는 타입을 뜻한다. %r은 32비트 정수, %rd는 64비트 정수, %f는 32비트 float, %p는 1비트 predicate이다..
PTX는 예상보다 더 ‘장황한 표기’에 가깝다. 예를 들어 %rd6에 주소 하나를 만드는 데 PTX 명령 세 개가 든다. 이는 PTX가 디바이스 불가지론적이기 때문에 생긴다.
왜 세 개인가?
CUDA 포인터는 기본적으로 “generic”이어서 global, shared, local 메모리를 가리킬 수 있다. cvta.to.global은 이 포인터가 global window 안에 있음을 명시하므로, 뒤에서 더 저렴한 ld.global을 쓸 수 있다. 그다음 mul.wide.s32는 인덱스 i에 4(sizeof(float))를 곱해 바이트 오프셋으로 만들고, 동시에 32→64비트로 확장한다. add.s64는 그것을 베이스 포인터에 더한다.
다음으로 ptxas는 디바이스 불가지론적인 PTX를, 그렇지 않은 아키텍처별 SASS로 변환한다. 생성된 SASS는 꽤 다르게 보인다.
$ cuobjdump -sass vadd
/*0000*/ MOV R1, c[0x0][0x28] ; // set up the stack pointer (ABI; unused here)
/*0010*/ S2R R6, SR_CTAID.X ; // R6 = blockIdx.x
/*0020*/ S2R R3, SR_TID.X ; // R3 = threadIdx.x
/*0030*/ IMAD R6, R6, c[0x0][0x0], R3 ; // i = ctaid*ntid + tid
/*0040*/ ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ;// P0 = (i >= n)
/*0050*/ @P0 EXIT ; // if so, exit
/*0060*/ MOV R7, 0x4 ; // load literal 4 (sizeof(float)) into R7 as multiplier
/*0070*/ ULDC.64 UR4, c[0x0][0x118] ; // uniform load of a driver-provided system value
/*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i]
/*0090*/ IMAD.WIDE R2, R6, R7, c[0x0][0x160] ; // &a[i]
/*00a0*/ LDG.E R4, [R4.64] ; // b[i]
/*00b0*/ LDG.E R3, [R2.64] ; // a[i]
/*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ; // &c[i]
/*00d0*/ FADD R9, R4, R3 ; // a[i] + b[i]
/*00e0*/ STG.E [R6.64], R9 ; // c[i] = ...
/*00f0*/ EXIT ;
S2R 줄이 하는 일
S2R은 “special register to register”이다. 하드웨어가 스레드마다 유지하는 special 레지스터를 일반 레지스터로 복사해 IMAD가 산술에 사용할 수 있게 한다. 여기서는 SR_CTAID.X(블록의 인덱스, blockIdx.x)와 SR_TID.X(블록 내 lane의 인덱스, threadIdx.x)가 해당한다.
열 개 남짓한 가상 레지스터가 일곱 개의 실제 레지스터로 줄어들었다. ncu는 launch__registers_per_thread = 16이라고 보고한다. 디스어셈블리는 R9까지만 이름을 붙이지만, 할당기는 ABI와 정렬을 위해 몇 개를 더 예약한다.. 두 개의 mul.wide와 add 시퀀스는 하나의 IMAD.WIDE로 융합되었다. cvta 변환은 사라지고 주소 계산에 흡수되었다.
c[0x0][…] 피연산자는 constant bank 0이다. 드라이버가 관리하는 작은 영역으로, 여기에는 커널의 인자들—포인터 a, b, c와 크기 n—그리고 launch geometry가 들어 있다. 이 bank를 채우는 일은 QMD라는 구조체가 담당하며, 드라이버가 launch 시점에 GPU에 넘긴다. launch 자체가 카드에 도달하는 부분에서 다시 보게 된다.
왜 인자들이 constant bank 0에 있고, 어디에 있는가
이 인자들이 constant memory에 있는 이유는 이것이 broadcast 읽기이기 때문이다. 그리드의 모든 스레드가 동일한 포인터들을 필요로 하고, constant cache는 32개 lane 전체에 한 번에 공급할 수 있다. 레이아웃은 고정되어 있다. 0x160, 0x168, 0x170은 포인터 a, b, c이고, 0x178은 n이다. launch geometry는 0x0(blockDim.x) 근처에 있다. Bank 0에는 ABI 파라미터도 들어 있는데, 예를 들어 c[0x0][0x28]은 진입 시 MOV R1, c[0x0][0x28]가 읽는 스택 베이스다. 호스트 스텁이 launch용 인자를 패킹할 때도 같은 오프셋들이 다시 등장한다.
이 SASS를 담은 ‘cubin’ 파일은 ELF 파일이다. 즉, Linux가 일반 실행 파일과 공유 라이브러리에 사용하는 것과 같은 오브젝트 파일 컨테이너다. cuobjdump -elf를 보면 심볼 테이블, 기계어 코드를 담은 .text.vadd 섹션, 그리고 .nv.callgraph 같은 CUDA 전용 섹션을 볼 수 있다..
fatbinary 실행 파일은 cubin과 PTX를 하나의 ‘fatbin’으로 함께 묶고, 결과에 대해 cuobjdump를 실행해 보면 실행 파일 안에 둘 다 들어 있음을 알 수 있다.
$ cuobjdump vadd
...
Fatbin elf code: arch = sm_89 # the SASS we just read
Fatbin ptx code: arch = sm_89 compressed # the PTX, shipped too
이 4090에서 실제로 실행되는 것은 SASS지만, PTX도 앞으로의 호환성을 위한 fallback으로 함께 실린다. 이 바이너리를 cubin이 지원하지 않는 아키텍처의 GPU로 가져가면, 드라이버는 로드 시점에 PTX를 새 SASS로 JIT 컴파일할 수 있다.
마지막으로, 이 fatbin은 호스트 실행 파일 안에 중첩되어 들어가며, readelf -S를 사용하면 자체 섹션을 차지하고 있음을 볼 수 있다.
$ readelf -S vadd
...
[18] .nv_fatbin PROGBITS ...
[19] __nv_module_id PROGBITS ...
[29] .nvFatBinSegment PROGBITS ...
...
nvcc가 뱉는 vadd 바이너리는 호스트 코드, Ada SASS를 담은 완전한 ELF 오브젝트, 그리고 PTX 사본을 포함한 단일 실행 파일이다. PTX는 장황한 평문 텍스트이므로, nvcc는 바이너리 크기를 작게 유지하기 위해 기본적으로 이를 압축한다. 드라이버는 미리 컴파일된 SASS가 지원하지 않는 아키텍처에서 바이너리가 실행될 때만 그것을 압축 해제하고 JIT 컴파일한다.
컴파일된 GPU 기계어는 이제 ./vadd 실행 파일의 .nv_fatbin 섹션 안에 가만히 들어 있다. 호스트에서 프로그램을 실행할 때는 두 세계를 이어야 한다. 호스트 CPU와 PCIe 버스 건너편에 있는 GPU다.
다리를 건너는 법을 아는 호스트 바이너리를 준비하기 위해 프런트엔드 컴파일러(cudafe++)는 코드 안에 숨겨진 constructor를 삽입한다. 이것은 main 함수가 시작되기 전에 실행된다. 그 역할은 내장된 fatbinary를 CUDA 런타임에 등록하고, 나중에 런타임이 사용할 매핑을 기록하는 것이다. 즉, 호스트 쪽 함수 포인터 vadd를 fatbin 안에 있는 컴파일된 디바이스 커널의 mangled name과 연결한다.
컴파일러가 vadd<<<4096, 256>>>(da, db, dc, n)를 만나면, 그 고수준 표현식을 생성된 호스트 launch stub으로 치환한다. 이 스텁은 커널 인자들을 호스트 메모리의 버퍼에 패킹한다. 포인터 da, db, dc와 정수 n은 각각 바이트 오프셋 0, 8, 16, 24에 정렬된다. 이 오프셋들은 앞서 SASS 기계어가 constant bank 0에서 읽는 것을 보았던 0x160, 0x168, 0x170, 0x178 오프셋들에 대응한다..
// from vadd.cudafe1.stub.c
void __device_stub__Z4vaddPKfS0_Pfi(const float *__par0, const float *__par1,
float *__par2, int __par3) {
__cudaLaunchPrologue(4);
__cudaSetupArgSimple(__par0, 0UL); // arg buffer offset 0
__cudaSetupArgSimple(__par1, 8UL); // offset 8
__cudaSetupArgSimple(__par2, 16UL); // offset 16
__cudaSetupArgSimple(__par3, 24UL); // offset 24
__cudaLaunch((char*)(void(*)(const float*, const float*, float*, int))vadd);
}
인자가 패킹되면 스텁은 __cudaLaunch를 호출하고, 호스트 쪽 더미 vadd 함수의 메모리 주소를 넘긴다. 이 호스트 함수는 CPU에서는 빈 껍데기에 불과하므로, 그 호스트 메모리 주소는 조회 키 역할을 한다. 런타임은 이 주소로 등록 테이블을 조회해 대응하는 디바이스 쪽 심볼 이름을 찾고, 이어서 닫힌 소스의 사용자 모드 드라이버(libcuda.so.1)로 경계를 넘어가 그 커널 launch를 시작한다. 드라이버의 사용자 모드 부분은 CUDA 툴킷이 아니라 GPU의 커널 드라이버와 함께 제공된다. strace에서 보이는 libcuda.so.1은 이 시스템의 드라이버 릴리스인 libcuda.so.590.48.01로 해석된다..
런타임은 프로그램의 첫 GPU 호출에서 이 드라이버를 동적으로 연다. strace로 이를 잡아낼 수 있다.
$ strace -f -e trace=openat ./vadd
...
openat(..., "/lib/x86_64-linux-gnu/libcuda.so.1", O_RDONLY|O_CLOEXEC) = 3
...
이 첫 호출이 수행될 때 ‘context’가 생성된다. 여기에는 드라이버가 디바이스와 통신하는 데 필요한 모든 인프라가 들어 있으며, 그중에는 CPU가 GPU와 대화하는 _channel_도 포함된다. 이에 대해서는 다음 섹션에서 더 이야기하겠다.
이 단계에서도 컴파일된 기계어는 아직 GPU에 도달하지 않았다. CUDA 12.2 이후로는 모듈 로딩이 기본적으로 lazy다. CUDA_MODULE_LOADING으로 제어된다. CUDA 11.7에서는 opt-in으로 도입되었고 오랫동안 기본값은 EAGER였지만, 12.x 계열에서 기본값이 LAZY로 바뀌었다. 원한다면 로딩 비용을 미리 지불하도록 덮어쓸 수 있다.—즉, 드라이버는 특정 커널이 실제로 처음 launch될 때까지 그 커널의 SASS cubin을 카드 메모리에 올리는 일을 미룬다.
libcuda 아래에는 커널 모드 드라이버 nvidia.ko가 있고, libcuda는 디바이스 파일에 ioctl을 호출해 여기에 접근한다. cuLaunchKernel이 마침내 GPU에 작업을 올려야 할 때, 그것은 이 커널 모듈과의 대화가 된다. 이어지는 내용은 그 대화의 기계적 세부다.
GPU는 CPU처럼 함수 호출을 받지 않는다. 점프할 entry point도 없고, CPU에서 인자를 밀어 넣을 스택도 없다. GPU는 PCIe 버스 너머에 앉아 호스트 메모리에 있는 드라이버 명령 스트림을 읽는다. 이 지점 이후 cuLaunchKernel이 하는 모든 일은 완성된 launch 명령 하나를 그 스트림에 넣고, GPU에 그것이 준비되었다고 알리기 위한 것이다.
가장 먼저 해야 할 일은 GPU 코드를 디바이스에 적재하는 것이다. vadd를 처음 실행할 때 드라이버는 커널 코드를 복사해 넣는다. 버퍼를 할당하고 SASS를 복사한다.
코드가 GPU에 올라가면, CPU는 GPU가 그것을 읽고 실행을 시작하도록 만들어야 한다. 이것은 호스트 메모리와 디바이스 메모리를 오가는 복잡한 춤을 통해 이뤄진다. 호스트와 GPU는 서로의 메모리 공간 일부를 매핑할 수 있지만, PCIe 버스를 가로지르는 접근에는 비용이 든다. 커널 launch를 달성하려면 양쪽 모두 서로 다른 공간에 사는 여러 구조체에 써야 한다. 이 구조체들이 바로 channel—GPU의 작업 큐—를 이룬다.
호스트 RAM에 사는 중요한 구조체는 두 개다. pushbuffer와 GPFIFO이며, 둘이 합쳐 GPU가 수행해야 할 작업 목록을 나타낸다.
pushbuffer는 드라이버가 GPU에 대한 명령, 즉 _method_를 쓰는 메모리 영역이다. method는 GPU 고유의 명령 인코딩으로 표현된 레지스터 주소와 값의 쌍이며, 이 쌍이 GPU가 수행할 동작을 정의한다.
GPFIFO는 포인터들로 이루어진 링 버퍼로, GPU와 CPU가 GPU가 아직 읽어야 할 것과 이미 읽은 것을 조정하는 데 사용한다. GPFIFO의 각 엔트리는 두 개의 32비트 워드로 이루어지며, pushbuffer의 한 구간을 설명한다. 이 경우 base는 호스트 메모리를 가리키는 GPU 가상 주소다. (base, length).
GPU는 계속 GPFIFO를 따라가며 작업을 찾는다. 드라이버와 GPU 사이에는 두 개의 커서가 유지되어야 한다. GP_GET(GPU가 어디까지 소비했는가)와 GP_PUT(드라이버가 어디까지 생산했는가)다. 두 커서는 USERD라는 채널별 작은 구조체 안에 있으며, 여기서는 디바이스 메모리에 있다. 커널을 launch하려면 드라이버는 관련 method들로 pushbuffer 구간을 채우고, GPFIFO 엔트리가 그 구간을 가리키게 한 뒤 GP_PUT를 전진시킨다. GPU가 그 엔트리를 소비하면 GP_GET를 전진시킨다.
각 조각이 어디에 사는가.
우리의 launch는 먼저 SET_INLINE_QMD_ADDRESS_A/B로 시작하는 method 폭주에 의해 트리거되고, 그 뒤를 LOAD_INLINE_QMD_DATA 연속이 잇는다. 이것이 왜 이 method인지, libcuda가 닫힌 소스인데 어떻게 아는지는 부록를 보라.. 이 method들은 “Queue Meta Data”(QMD)라는 객체를 pushbuffer로 스트리밍하는 역할을 한다.
QMD는 compute grid의 launch descriptor다. 여기에는 그리드와 블록 차원—.cu 코드의 4096과 256—스레드당 레지스터 수와 필요한 shared memory, 그리고 두 개의 주소가 들어 있다. 하나는 프로그램의 시작 주소(첫 launch에서 GPU 메모리에 적재된 SASS)이고, 다른 하나는 커널 인자를 담은 constant bank의 주소다. 호스트 스텁이 패킹한 인자들은 여기에 들어간다. 드라이버가 그것들을 복사해 넣고 bank의 주소를 QMD에 기록한다. QMD는 GPU에게 SASS가 어디 있는지, 그 SASS를 어떻게 병렬 프로그램으로 바꿔야 하는지, 그리고 그 프로그램의 완료를 어디에 신호해야 하는지를 알려준다.
이제 GPU가 실행을 시작할 준비는 끝났다. 문제는 GPU의 host engine호스트와 인터페이스하는 GPU 제어 로직의 일부.이 아직 행동하지 않았다는 점이다. 최신 카드에서는 이 엔진이 커서를 감시하지 않는다. 예전 GPU는 USERD를 snoop해서 GP_PUT를 쓰는 것만으로 충분했지만, Turing 이후에는 그렇지 않으므로 드라이버가 대신 도어벨을 울린다.. 그래서 GP_PUT의 변경은 엔진에게 보라고 알려줄 무언가가 있기 전까지 그냥 거기 머무른다.
그 엔진에게 보라고 알리는 것이 doorbell이다. GPU는 자신의 레지스터 일부를 프로세스에 작은 윈도로 매핑하고, 그중 하나가 도어벨이다. 드라이버는 여기에 채널의 _work-submit token_을 쓴다. 이 토큰은 어느 채널에 새 작업이 생겼는지 알려준다.
도어벨이 울리면 host engine은 갱신된 GP_PUT를 읽고, 새로운 GPFIFO 엔트리를 따라가 pushbuffer 구간으로 간 뒤, DMA로 그 안의 method들을 가져온다. 그리고 우리의 QMD를 담은 compute method에 도달하면, 그 descriptor를 “compute work distributor”에 넘긴다. 이에 대해서는 곧 더 설명하겠다.
CPU 쪽에서 보면 launch는 끝났다. cuLaunchKernel은 도어벨이 울리는 순간 반환된다. 호출은 비동기이므로 제어는 프로그램으로 돌아오고 CPU는 GPU가 일하는 동안 계속 실행된다. 커널이 실행된 뒤 호스트 쪽 이야기를 다시 이어받는다.
이제 GPU가 자기 일을 시작할 차례다.
host engine은 QMD를 compute work distributor에 넘긴다. 여전히 GigaThread Engine이라고 부르기도 한다. GPU 전체에 하나만 있다.. VRAM에는 하나의 선형 SASS 명령 목록이 있고, compute work distributor와 QMD는 그 선형 스레드 명령 목록을 전체 Streaming Multiprocessor(SM)에 걸친 대규모 병렬 프로그램으로 만드는 과정을 하드웨어에 알려주는 첫 단계다.
스택을 내려온 우리의 여정에서, 이제 compute work distributor는 256개 스레드의 블록 4096개를 설명하는 QMD를 갖고 있다. 우리가 대상으로 하는 카드는 128개 SM을 가진 GeForce RTX 4090 칩이다. NVIDIA의 AD102-300-A1 SKU는 전체 다이의 물리적 144개 SM 중 16개를 비활성화해 제조 수율을 높인다. 자세한 내용은 NVIDIA Ada GPU Architecture whitepaper를 보라.. distributor의 임무는 이 128개를 모두 작업으로 가득 채워 두는 것이다.
컴파일된 기계어는 global memory에 단일한 선형 시퀀스로 놓여 있다. 각 SM은 자체 로컬 Instruction Cache(I-cache)를 가지고, GPU의 각 활성 워프는 각자의 사적인 Program Counter(PC)를 유지한다. Volta 이후에는 이 모델이 더 세분화되어 각 _스레드_가 자신만의 프로그램 카운터와 호출 스택을 갖는다(Independent Thread Scheduling). 그래서 워프 안의 스레드가 자유롭게 divergence와 reconvergence를 할 수 있다. 하지만 issue는 여전히 워프 단위다. 매 사이클마다 스케줄러는 워프 하나를 고르고 현재 공통 PC에 있는 lane들에 issue한다.. 그러면 SM의 스케줄러들이 그 선형 시퀀스에서 독립적으로 명령을 fetch할 수 있어, 서로 다른 워프가 같은 SASS 코드를 다른 속도로 실행하거나 서로 다른 분기 경로를 따라갈 수 있다.
VRAM에는 하나의 명령 스트림이 있고, 각 SM에서 로컬로 캐시된다. SM은 최대 48개의 resident warp를 유지하지만(그리드), 네 개의 스케줄러는 사이클당 최대 한 개씩만 명령을 issue한다. 여기서는 거의 모든 워프가 LDG.E load(주황색)에 걸려 멈춰 있고, FADD(초록색)를 issue하는 슬롯은 하나뿐이다.
우리 SM들의 하드웨어 제약은 동시에 실행될 수 있는 블록 수를 정한다. cudaGetDeviceProperties가 이 정보를 알려준다.
+------------------------------------------------------------+
| AD102 SM Resource Caps |
+------------------------------------------------------------+
| Max Active Threads/SM | 1,536 threads (48 warps) |
| Register File/SM | 65,536 32-bit registers (256 KB) |
| Shared Memory/SM | 100 KB |
+------------------------------------------------------------+
우리의 launch configuration은 256개 스레드(8 warps) 블록을 지정하고, ptxas는 스레드당 16개 레지스터를 예약했다.
스레드 수용량이 더 빡빡한 병목이므로, 각 SM에는 동시에 최대 **6개 블록(48 warps)**만 올라갈 수 있다.
distributor는 이 6개의 resident block을 SM 하나에 배정한다. 각 SM은 **네 개의 processing block(sub-partition)**으로 나뉜다. 각 sub-partition은 자체적으로 완결된 실행 파이프라인이다.
SM은 48개의 resident warp를 이 네 개의 sub-partition에 고르게 분배한다. 그래서 SM이 가득 찼을 때 각 warp scheduler는 관리할 12개의 활성 warp()를 갖게 된다. 매 사이클마다 warp scheduler는 이 12개 후보를 평가하고, _eligible_한 워프 하나를 골라 그 다음 명령을 자신의 실행 슬라이스에 있는 32개의 물리적 lane 전체로 디스패치한다.
GPU가 명령이 실행 준비가 되었는지를 판단하는 방식은 CPU와 다르다. 현대의 out-of-order CPU는 런타임에 동적으로 의존성을 찾아내며, reorder buffer와 rename logic에 실리콘을 써서 단일 스레드에서 병렬성을 뽑아낸다. GPU는 그럴 필요가 없다. 많은 워프를 resident 상태로 두고, 어떤 워프가 stall되면 다른 워프로 전환함으로써 지연을 숨긴다. 병렬성이 핵심인 환경에서, 무거운 의존성 추적 장치는 실리콘을 쓰는 올바른 방법이 아니다. 그래서 하드웨어는 타이밍을 예측할 수 있는 모든 것은 컴파일러가 스케줄링하도록 맡기고, 예측할 수 없는 것만 가벼운 하드웨어 scoreboard에 의존한다.
모든 128비트 SASS 명령에는 ptxas가 기록한 압축된 control-code payload가 실린다. 가장 명확한 공개 재구성은 Citadel 마이크로벤치마킹 논문(Jia et al., “Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking”)과 Maxwell용 이 maxas control-code 노트이다.. 이 스케줄링 control bit들은 하드웨어 타이밍을 직접 지시하며, 세 가지 핵심 지시를 담는다.
LDG)와 special function(MUFU)—에 대해 하드웨어는 워프마다 **여섯 개의 물리적 scoreboard barrier(0~5번)**를 제공한다.왜 디스어셈블리에서 이 비트들이 보이지 않는가
NVIDIA의 표준 nvdisasm 도구로 바이너리를 디스어셈블하면, 기본적으로 원시 control code는 숨겨진다. 도구가 그것들을 제거하고 깔끔한 표준 SASS 니모닉만 보여주기 때문이다. 하지만 control code는 실제로 명령 바로 옆에 저장되어 있다. cuobjdump -sass로 원시 바이너리를 보고 16진수 명령 주석(/* 0x... */ 등)을 자세히 보면, 이 control bit들을 담고 있는 압축된 원시 16진수 워드를 볼 수 있다.
정확한 레이아웃에 대해 우리가 아는 것은 마이크로벤치마킹 커뮤니티의 역공학 노력에서 나온다. 비트 필드는 Maxwell, Volta, Ampere, Ada Lovelace 사이에서 이동하고 진화했지만, 핵심 아키텍처 개념은 동일하다. 즉, 컴파일 시점의 정적 스케줄링 메타데이터가 명령 스트림 안에 직접 패킹되어, SM 하드웨어를 가능한 한 단순하고 전력 효율적으로 유지한다.
vadd에 대해 cuobjdump -sass를 실행하면, 각 명령에는 두 개의 64비트 워드로 된 원시 128비트 인코딩이 붙어 있고, 각 쌍의 두 번째 워드가 control payload를 담는다.
$ cuobjdump -sass vadd # control payload
/*00a0*/ LDG.E R4, [R4.64] /* 0x000ea8000c1e1900 */
/*00b0*/ LDG.E R3, [R2.64] /* 0x000ea2000c1e1900 */
/*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] /* 0x000fe200078e0207 */
/*00d0*/ FADD R9, R4, R3 /* 0x004fca0000000000 */
/*00e0*/ STG.E [R6.64], R9 /* 0x000fe2000c101904 */
control payload를 추출해 보면—가장 명확한 공개 재구성은 Citadel 마이크로벤치마킹 논문(Jia et al., “Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking”)과 Maxwell용 이 maxas control-code 노트이고, 비트 레이아웃은 부록에 있다—ptxas가 기록한 스케줄을 각 지시가 실제로 작동하는 모습과 함께 볼 수 있다.
| instruction | stall | yield | sets | waits-on |
|---|---|---|---|---|
LDG.E | 4 | yes | B2 | — |
LDG.E | 1 | yes | B2 | — |
IMAD.WIDE | 1 | yes | — | — |
FADD | 5 | no | — | B2 |
STG.E | 1 | yes | — | — |
두 load는 지시 3을 활용해 각각 같은 scoreboard barrier인 B2를 “set”한다. 그리고 load된 R4와 R3가 처음 필요한 명령인 FADD는 B2를 기다리도록 되어 있다. 두 load가 모두 돌아와 barrier가 해제될 때까지 이 워프는 ineligible 상태이며, 스케줄러는 이 워프를 건너뛰고 같은 sub-partition의 다른 열한 개 워프 중 하나를 고른다.
FADD→STG의 인계는 지시 1의 사례다. 부동소수점 덧셈은 지연이 고정되어 있으므로 barrier가 필요 없다. 대신 FADD는 stall=5를 들고 있으며, STG가 R9를 읽기 전에 그 값이 도착하는 데 필요한 몇 사이클 동안 워프를 멈춰 세운다.
지시 2인 yield bit는 이 시퀀스 전반에서 켜졌다 꺼졌다 하며, 컴파일러가 곧 대기할 연산 주변에서 스케줄링 우선순위를 미세 조정한다.
매 사이클마다 스케줄러는 워프의 6비트 barrier 상태와 작은 stall counter를 읽고, 각 워프에 대한 eligibility 결정을 내린다. 이것이 GPU가 거의 0에 가까운 하드웨어 스케줄링 오버헤드로 지연을 숨기는 방식이다.
warp scheduler가 eligible한 워프를 찾아 LDG.E load를 issue하면, 그 하드웨어 요청이 메모리 계층을 내려가는 경로를 따라갈 수 있다. 워프의 32개 스레드는 각각 주소를 계산한다. 우리의 스레드들은 float 배열의 연속된 원소들(각각 4바이트)에 접근하므로, 워프는 연속된 128바이트 블록( bytes)을 요청한다.
SM의 load/store unit는 이 연속 접근 패턴을 감지해 request coalescing을 수행한다. 스레드별 4바이트 요청 32개를 32바이트 sector 요청 네 개로 병합한다. fetch 단위가 32바이트이므로 이상적이다. 만약 읽기가 이렇게 연속적이고 coalesced되지 않았다면 필요한 것보다 더 많은 데이터를 읽어오게 된다.
병합된 요청은 먼저 SM의 로컬 L1 Data Cache를 확인한다. 거기서 miss가 나면, 모든 128개 SM을 분산된 72 MB L2 Cache 슬라이스들과 연결하는 고대역폭 crossbar interconnect를 통해 라우팅된다. 요청이 L2 cache에서도 miss 나면 더 아래의 memory controller로 내려가고, 메모리 버스를 건너 물리적인 GDDR6X VRAM 칩으로 이동한다. RTX 4090은 A100이나 H100 같은 데이터센터급 GPU에서 쓰이는 High-Bandwidth Memory(HBM)가 아니라 GDDR6X 메모리를 사용한다.. 루프 마지막에 c[i]를 쓰는 STG.E store도 원칙적으로는 정확히 같은 경로를 반대로 따른다. 다만 뒤에서 보겠지만 c[i]는 실제로 VRAM까지는 가지 않는다..
컴파일한 커널을 NVIDIA Nsight Compute 프로파일러(ncu) 아래서 실행하면, 의미 있는 메트릭 몇 가지를 얻을 수 있다.
$ ncu --metrics \
launch__grid_size,launch__block_size,launch__registers_per_thread,\
launch__waves_per_multiprocessor,sm__warps_active.avg.pct_of_peak,\
smsp__issue_active.avg.pct_of_peak,dram__throughput.avg.pct_of_peak,\
gpu__time_duration.sum \
./vadd
...
----------------------------------------------------------
Metric Name Unit Value
----------------------------------------------------------
launch__grid_size 4,096
launch__block_size 256
launch__registers_per_thread 16
launch__waves_per_multiprocessor 5.33
sm__warps_active.avg.pct_of_peak % 82.77
smsp__issue_active.avg.pct_of_peak % 5.17
dram__throughput.avg.pct_of_peak % 79.65
gpu__time_duration.sum us 10.78
----------------------------------------------------------
실행 동안 워프의 82.77%가 active 상태였다. 워프가 실제로 명령을 issue한 시간은 5.17%였다. DRAM은 최대 활용도의 79.65%로 동작하고 있었다.
이 커널은 arithmetic intensity가 극히 낮다. 12바이트의 데이터를 전송할 때마다(4바이트 load 두 번과 4바이트 store 한 번) 정확히 한 번의 부동소수점 덧셈(FADD)과 약간의 포인터 산술만 수행한다.
따라서 10.78s는 결국 DRAM 버스가 얼마나 빠르게 커널에 입력을 공급할 수 있느냐로 귀결된다. 여기서는 최대치의 약 5분의 4 수준이다. 버스를 건너는 것은 전체 12 MB가 아니라 두 입력뿐이다. ncu를 보면 DRAM에서 8.4 MB를 읽고 쓴 것은 사실상 없다. 4 MB 출력 c는 72 MB L2에 들어맞아서, 나중의 device-to-host copy가 다시 읽어가기 전까지 DRAM으로 flush되지 않는다. 최대치의 5분의 4라는 수치는 읽기 쪽 기준이다. 8.4 MB / 10.78 s 780 GB/s..
결과는 이제 GPU의 L2 cache 안에 있다. 터미널을 구동하는 것은 CPU이므로, 우리에게 보여주려면 CPU가 이 결과를 가져와야 한다. 다시 CPU 관점의 사건으로 돌아가자.
launch는 도어벨이 울린 순간 CPU에 제어를 돌려주었다. 그러므로 GPU는 CPU에 자신이 끝났다고 알려야 한다. 4096개 블록 중 마지막 블록이 retire될 때, GPU는 QMD가 들고 있던 completion semaphore를 게시함으로써 그렇게 한다. fence field는 워드 23–24에 있다..
device-to-host cudaMemcpy(c, dc, …)고정 메모리의 cudaMemcpyAsync라면 대기를 건너뛰고 호스트가 앞서 나가게 할 수 있다.는 default stream에서 커널 뒤에 놓여 있으므로, 이 전송을 수행하는 GPU의 copy engine은 이 semaphore에 의해 게이트된다. 값이 나타나면 GPU는 DMA를 수행한다. c는 여전히 72 MB L2 안에 dirty 상태로 남아 있으므로—STG.E store가 DRAM으로 spill될 필요가 전혀 없었기 때문에—engine의 읽기는 곧바로 L2에서 서비스되고, 데이터는 DRAM 왕복 없이 PCIe를 건넌다.
copy가 끝나면 그것은 자기 자신의 semaphore를 게시하고, 호스트는 cudaMemcpy 안에서 그것을 기다리고 있었다. 호스트에서 cudaMemcpy가 완료되면 c는 다시 일반적인 호스트 메모리가 되고, printf는 RAM에서 c[0]과 c[n-1]를 읽어 문자열로 포맷한 뒤 stdout에 대한 write syscall로 넘긴다.
커널 소스는 cicc를 거쳐 PTX가 되었고, ptxas를 거쳐 SASS가 되었다. fatbinary는 이것을 PTX의 fallback 사본과 함께 cubin을 품은 fatbin으로 패킹했고, 링커는 그것을 평범한 Linux 실행 파일 안에 용접해 넣었다. constructor는 main 전에 그 fatbin을 등록해 호스트 스텁을 mangled device name에 매핑했다. 첫 launch는 lazily cubin을 GPU에 업로드했다. cuLaunchKernel은 launch configuration으로부터 QMD를 만들고, 그것을 GPU method로 pushbuffer에 써 넣고, GP_PUT를 전진시킨 뒤 단일 MMIO store로 도어벨을 울렸다. 그러자 GPU의 host engine이 작업을 가져와 QMD를 compute work distributor에 넘겼다. distributor는 4096개 블록을 128개 SM에 완전 점유 상태로 퍼뜨렸고, SM당 네 개의 warp scheduler가 컴파일러가 써 둔 stall count를 포함한 128비트 명령을 issue했다. 그리고 coalesced된 메모리 경로가 입력을 DRAM에서 최대 대역폭의 5분의 4 속도로 끌어와, 백만 개 lane 각각에서 단 하나의 합을 계산했다. 완료 semaphore와 copy engine이 그 결과를 다시 버스를 건너 printf가 기다리는 곳으로 가져왔고, 우리는 다음 사실을 알게 되었다.
c[0]=2.000000 c[n-1]=2.000000
Claude와 나는 여기서 커널 launch의 각 부분이 일어나는 모습을 보기 위해 아주 다양한 요령을 사용했다. 그중 일부는 오픈 커널 모듈을 끈질기게 읽어 얻은 것이다.
이 글의 몇몇 주장은 libcuda가 닫힌 소스이기 때문에 오픈 소스만으로는 읽어낼 수 없다. 그것들을 알아내기 위해 쓸 만한 진단용 훅이 몇 가지 있다.
드라이버의 method write는 syscall을 거치지 않는다(드라이버가 이미 매핑해 둔 write-combined buffer에 직접 쓰기 때문이다). 그래서 그것을 찾으려면 메모리를 읽어야 한다. 우리는 mmap을 감싸는 LD_PRELOAD shim을 사용했다. 이것은 드라이버가 /dev/nvidia* 파일에서 매핑하는 모든 영역을 기록하고, launch가 반환된 직후 테스트 프로그램이 호출하는 함수를 노출해 그 영역들을 dump한다.
#define _GNU_SOURCE
#include <stdio.h>
#include <stdlib.h>
#include <dlfcn.h>
#include <sys/mman.h>
#include <unistd.h>
#include <string.h>
// Dynamic linker function pointers
static void* (*orig_mmap)(void*, size_t, int, int, int, off_t) = NULL;
// Store captured channel mappings
struct Map {
void* addr;
size_t length;
off_t offset;
char path[256];
} maps[128];
static int map_count = 0;
void* mmap(void* addr, size_t length, int prot, int flags, int fd, off_t offset) {
if (!orig_mmap) {
orig_mmap = dlsym(RTLD_NEXT, "mmap");
}
void* ret = orig_mmap(addr, length, prot, flags, fd, offset);
if (ret != MAP_FAILED && fd != -1 && map_count < 128) {
char proclink[256];
char path[256];
sprintf(proclink, "/proc/self/fd/%d", fd);
ssize_t len = readlink(proclink, path, sizeof(path) - 1);
if (len != -1) {
path[len] = '\0';
// We care about NVIDIA device files
if (strstr(path, "/dev/nvidia")) {
maps[map_count].addr = ret;
maps[map_count].length = length;
maps[map_count].offset = offset;
strcpy(maps[map_count].path, path);
map_count++;
}
}
}
return ret;
}
// Expose a function to dump memory ranges holding the pushbuffer
void dump_pushbuffer() {
printf("\n=== [Shim] Dump of Mapped Pushbuffers ===\n");
for (int i = 0; i < map_count; i++) {
// User-space channels/pushbuffers are mapped at large sizes
if (maps[i].length >= 0x1000) {
unsigned int* ptr = (unsigned int*)maps[i].addr;
printf("Mapping %d: %s, at %p (%zu bytes), offset 0x%lx\n",
i, maps[i].path, maps[i].addr, maps[i].length, (long)maps[i].offset);
// Walk the words looking for a method-header burst
for (size_t j = 0; j < maps[i].length / 4; j++) {
unsigned int word = ptr[j];
unsigned int opcode = (word >> 29) & 0x7; // 1 = INC
unsigned int count = (word >> 16) & 0x1FFF; // payload words
unsigned int method = (word & 0xFFF) << 2; // register offset
// 0x318 is SET_INLINE_QMD_ADDRESS_A, the start of the inline burst
if (opcode == 1 && method == 0x318) {
printf(" [+] Method burst at word %zu: header = 0x%08X\n", j, word);
printf(" INC, count %d, offset 0x%04X\n", count, method);
for (unsigned int k = 1; k <= count && (j + k) < (maps[i].length / 4); k++) {
printf(" word %02u: 0x%08X\n", k, ptr[j + k]);
}
}
}
}
}
}
이것을 공유 라이브러리로 컴파일한다.
$ gcc -shared -fPIC -o shim.so shim.c -ldl
그리고 테스트 프로그램에서 커널 launch 직후 dump_pushbuffer()를 호출하고, 이 mmap이 libc의 것 대신 실행되도록 shim을 preload한 상태로 실행한다.
$ LD_PRELOAD=./shim.so ./vadd
드라이버는 채널을 위해 write-combined buffer를 매핑하고, dump 코드는 그것을 순회하며 launch의 method burst를 출력한다. 이제 그것을 디코드해야 한다.
pushbuffer method는 헤더 워드 하나와 그 뒤의 데이터 워드들로 구성된다. 헤더는 네 개의 필드를 패킹한다(clc46f.h의 NVC46F_DMA_INCR_* 매크로로 정의됨).
0x1은 증가하는 method write(INC_METHOD/INCR_OPCODE_VALUE), 0x3은 비증가 method write(NON_INC_METHOD), 0x4는 즉시 데이터 write(IMMD_DATA_METHOD)다.NVC46F_DMA_INCR_COUNT).NVC46F_DMA_INCR_SUBCHANNEL).NVC46F_DMA_INCR_ADDRESS)이며, shim이 이를 다시 시프트해 원래 값으로 되돌린다.여기와 관련 있어 보이는 launch path는 두 개다. method들은 compute class별로 src/common/sdk/nvidia/inc/class/에 정의되어 있다 — clc3c0.h(Volta), clc5c0.h(Turing), clc6c0.h/clc7c0.h(Ampere), clc9c0.h(Ada), clcbc0.h(Hopper), clcdc0.h(Blackwell). Ada 헤더(clc9c0.h)는 클래스 번호 0xC9C0만 정의하고 Ampere method 집합을 상속하는 29줄짜리 스텁이므로, 우리가 실제로 읽는 정의는 Ampere 헤더에 있다.
0x0318 — SET_INLINE_QMD_ADDRESS_A(Ampere 헤더의 NVC6C0_SET_INLINE_QMD_ADDRESS_A로 정의되며, Ada에서도 변경 없이 상속됨). 이것은 LOAD_INLINE_QMD_DATA(i)(오프셋 0x0320 + i * 4)를 통해 QMD를 pushbuffer로 직접 스트리밍하는 inline-QMD burst를 연다.0x02b4 — SEND_PCAS_A, out-of-line 경로로, VRAM 어딘가에 있는 QMD에 대한 포인터만 전달한다.dump를 보면 pushbuffer에 실제로 어느 쪽이 들어가는지 알 수 있다. dump는 inline 경로를 보여준다. 하나의 increasing-method burst가 count 66으로 SET_INLINE_QMD_ADDRESS_A에서 열리고 있다. 그 66개 워드는 두 개의 주소 워드(SET_INLINE_QMD_ADDRESS_A/_B, 0x0318/0x031c)와 64개의 LOAD_INLINE_QMD_DATA 워드(0x0320부터)를 뜻한다. 즉, 256바이트짜리 QMD 하나가 inline으로 실린 것이다. 그 안에서 워드 12는 0x1000, 워드 18은 0x100인데, 이는 vadd<<<4096, 256>>>의 4096과 256이다.
Queue Meta Data(QMD) 구조체는 src/common/sdk/nvidia/inc/class/cla0c0qmd.h 안에서 32비트 경계를 가로지르는 다중 워드(MW) 비트 필드로 정의된 다중 워드 레이아웃으로 표현된다. QMD는 주소처럼 보이는 필드를 여러 개 저장하지만, 모두 같은 종류의 값은 아니다.
PROGRAM_OFFSET — MW(287:256) (Word 8)은 32비트 entry-point 오프셋이며, 채널의 코드 베이스에 대한 상대값이지 64비트 포인터가 아니다.CONSTANT_BUFFER_ADDR_LOWER(i) / ADDR_UPPER(i) — MW(959+i 64:928+i 64) (예를 들어 우리의 인자를 담고 있는 Constant Bank 0은 Words 29–30에 있다).RELEASE0_ADDRESS_LOWER/UPPER — MW(767:736) (Words 23–24), fence/semaphore에 사용된다.CIRCULAR_QUEUE_ADDR_LOWER/UPPER — MW(319:288) (Words 9-10).이 값들은 CPU가 직접 읽을 수 없는 디바이스 메모리를 가리킨다. 평범한 load는 fault를 내고, cudaMemcpy와 cuMemcpyDtoH도 이 주소를 거부한다.
그래서 GPU로 읽어야 한다. 아래의 작은 커널은 raw pointer에서 512바이트를 읽어 호스트가 가져올 수 있는 버퍼로 복사한다.
__global__ void peek(const unsigned char* src, unsigned char* dst) {
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < 512;
i += blockDim.x * gridDim.x) {
dst[i] = src[i];
}
}
QMD의 각 필드를 가리키도록 해 보면, 그중 정확히 하나만 512바이트 전부에 대해 SASS와 일치한다. QMD 안에서 유효한 GPU 가상 주소를 찾기 위해 메모리 스캐닝 shim을 실행해 보면, Word 48에서 일치 항목이 보인다.
qmd[48] -> 0x74167b272300 512 / 512 bytes match
왜 드라이버의 프로그램 필드는 Word 8의 PROGRAM_OFFSET인데, SASS는 Word 48(qmd[48])에서 일치하는가?
Word 8은 드라이버가 설정하는 32비트 오프셋만 담고 있는 반면, Words 48/49는 하드웨어 전용 HW_ONLY_INNER_GET(MW(1566:1536))과 HW_ONLY_INNER_PUT(MW(1598:1568)) 필드다. launch 이후 dump에서 이 워드들은 완전한 64비트 GPU 가상 주소를 담고 있고, Word 48의 값을 역참조하면 커널 SASS가 나온다. 가장 단순한 해석은 스케줄러가 launch 시점에 program offset을 이 스케줄러 전용 필드들로 해석해 넣는다는 것이다.
command stream은 메모리에서 읽어야 하지만, libcuda는 메모리와 GPU 객체들을 일반적인 방식으로 설정한다. 즉, 드라이버의 디바이스 파일에 ioctl을 수행한다(Michael Kerrisk, The Linux Programming Interface, Chapter 4 & 15 참고). 한 커널짜리 프로그램에 대해 strace를 돌리면 948개의 호출이 기록된다. 거의 모두 일회성 설정이고, 지속적인 launch 루프에서는 훨씬 적다., 대부분 두 개의 파일 디스크립터—/dev/nvidiactl과 /dev/nvidia-uvm—에 집중되어 있다.
$ strace -f -e trace=ioctl ./vadd
...
ioctl(8, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2a, 0x900), ...) # /dev/nvidiactl
ioctl(8, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2b, 0x30), ...) # /dev/nvidiactl
ioctl(9, ...) # /dev/nvidia-uvm
...
매직 바이트 0x46은 'F'로, NVIDIA resource manager ioctl의 magic 값이다. ‘magic’ 바이트는 모든 NVIDIA ioctl이 sanity check용으로 갖고 있는 값이다. Linux 커널 문서 참고.. command 번호는 오픈 커널 모듈의 nv_escape.h에 대조해 해석할 수 있다. 0x2A는 NV_ESC_RM_CONTROL이고 0x2B는 NV_ESC_RM_ALLOC이다.
eligibility 섹션에서 다룬 stall count, barrier, yield bit는 ptxas가 각 명령의 두 번째 64비트 워드 상단에 패킹하는 21비트 control field에서 나온다. cuobjdump -sass는 이를 니모닉 옆에 출력한다.
20 17 16 11 10 8 7 5 4 3 0
┌────────┬───────────┬──────┬──────┬─┬──────┐
│ reuse │ wait mask │ read │write │Y│stall │
│ (4) │ (6) │ barr │ barr │ │ (4) │
└────────┴───────────┴──────┴──────┴─┴──────┘
두 개의 3비트 인덱스는 명령이 설정하는 scoreboard barrier를 가리키고, 6비트 mask는 명령이 기다리는 barrier를 뜻하며, Y는 yield bit, stall은 정적 사이클 수다. 이 레이아웃은 문서화되어 있지 않으며, 마이크로벤치마킹으로 재구성된 것이다. 가장 명확한 공개 재구성은 Citadel 마이크로벤치마킹 논문(Jia et al., “Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking”)과 Maxwell용 이 maxas control-code 노트다..
시작 시점에 GPU 코드를 등록하기 위해 컴파일러가 실제로 어떤 코드를 생성하는지 보고 싶다면, nvcc --keep로 컴파일해서 vadd.cudafe1.stub.c를 살펴보면 된다.
프로세스 시작 시 등록은 자동 생성된 constructor가 담당한다.
// from vadd.cudafe1.stub.c
static void __sti____cudaRegisterAll(void) __attribute__((__constructor__));
static void __nv_cudaEntityRegisterCallback(void **__T4) {
__cudaRegisterEntry(__T4, (void(*)(const float*, const float*, float*, int))vadd,
_Z4vaddPKfS0_Pfi, -1);
}
static void __sti____cudaRegisterAll(void) {
__cudaRegisterBinary(__nv_cudaEntityRegisterCallback);
}
__attribute__((__constructor__)) 지시문은 링커에게 main이 시작되기 전에 __sti____cudaRegisterAll을 실행하라고 알려준다. 이것은 우리의 디바이스 바이너리를 CUDA 런타임에 등록하고 콜백을 예약한다. 실행되면 __cudaRegisterEntry는 호스트 함수 포인터 vadd를 mangled 디바이스 entry point _Z4vaddPKfS0_Pfi에 매핑해, cudaLaunchKernel이 launch 시점에 조회하는 해시 테이블을 만든다.