레지스터 압박과 스필

레지스터는 스레드가 사용할 수 있는 가장 빠른 저장 공간이지만, 각 SM이 보유한 수량에는 한계가 있습니다. 커널이 레지스터를 얼마나 많이 사용하는지를 레지스터 압박이라고 합니다.

상주하는 모든 스레드는 SM마다 하나씩 있는 레지스터 파일에서 공간을 할당받습니다. 스레드 하나에 필요한 레지스터가 많을수록 동시에 상주할 수 있는 스레드 수는 줄어듭니다.

레지스터 사용량이 높으면 SM에 들어갈 수 있는 워프 수가 직접 제한됩니다. 이렇게 점유율이 낮아지면 하드웨어가 지연 시간을 숨길 작업이 부족해질 수 있습니다.

스레드에 필요한 레지스터가 부족하면 컴파일러가 여분의 값을 밖으로 옮깁니다. 이 오버플로를 느린 메모리로 내보내는 레지스터 스필이라고 합니다.

스필된 값은 오프칩 DRAM에 있는 로컬 메모리로 이동합니다. 스필이 발생할 때마다 사용 가능한 레지스터 접근이 느린 왕복 메모리 접근으로 대체됩니다.

컴파일러에 사용량을 보고하도록 요청하십시오. nvcc에 --ptxas-options=-v를 추가하면 각 커널의 스레드당 레지스터 수와 스필 바이트가 출력됩니다.

보고서에는 스필 저장과 스필 읽기가 나열됩니다. 0이 아닌 스필 수는 커널이 느린 로컬 메모리 트래픽을 감당하고 있다는 경고 신호입니다.

컴파일할 때 상한을 설정할 수 있습니다. --maxrregcount 플래그는 스레드당 레지스터 수를 제한하여, 약간의 속도를 희생하는 대신 점유율을 높입니다.

커널별 힌트를 사용하는 편이 더 나은 경우가 많습니다. __launch_bounds__는 컴파일러에 블록 크기를 알려 주므로, 원하는 점유율에 맞춰 레지스터를 배분할 수 있습니다.

대개 한 번에 보유하는 값의 수를 줄이는 것만으로도 충분합니다. 계산 비용이 낮은 결과를 다시 계산하거나 변수의 범위를 좁히면 활성 상태로 유지되는 레지스터 수가 줄어듭니다.

ILP와 반복문 펼치기는 압박을 높이는 반면, 제한을 설정하면 점유율이 높아집니다. 가장 빠르게 실행되는 균형을 찾는 것이 중요하며, 그 지점을 알려 줄 수 있는 것은 프로파일러뿐입니다.

높은 레지스터 압박이 점유율을 낮추고 느린 DRAM으로 스필을 일으킬 수 있다는 것을 배웠습니다. 사용량을 측정하고 레지스터 수를 제한한 뒤 프로파일러의 안내를 따르십시오. 🎯