레지스터와 공유 메모리 한계

각 SM에는 고정된 양의 레지스터와 공유 메모리가 있습니다. 모든 상주 블록은 이 저장 공간에서 자신의 몫을 확보해야 합니다.

커널은 일정 수의 스레드당 레지스터를 사용합니다. 여기에 블록당 스레드 수를 곱하면 블록 하나에 필요한 레지스터 비용을 알 수 있습니다.

각 스레드에 많은 레지스터가 필요하면 들어갈 수 있는 스레드 수가 줄어들고, 따라서 상주 상태로 유지되는 블록 수도 줄어듭니다. 레지스터를 많이 사용하는 커널은 점유율이 낮아집니다.

이 플래그를 사용하여 컴파일하면 nvcc가 스레드당 레지스터 수를 출력하므로 커널이 레지스터를 과도하게 사용하는지 확인할 수 있습니다.

실행 제한을 사용하여 상한을 강제로 지정하면 컴파일러가 사용하는 레지스터 수를 줄이고 더 많은 워프를 상주 상태로 유지하게 할 수 있습니다.

각 블록은 공유 메모리를 요청할 수 있습니다. SM에는 공유 메모리 풀이 허용하는 수만큼의 블록만 들어가며, 보통 48~100KB입니다.

큰 __shared__ 타일을 요청하면 SM당 한두 개의 블록만 들어갑니다. 큰 타일은 데이터 재사용을 얻는 대신 점유율을 낮춥니다.

SM은 레지스터, 공유 메모리, 워프 한도를 기준으로 허용되는 블록 수를 계산합니다. 이 값들 중 가장 작은 값이 점유율을 결정합니다.

스레드에 허용된 것보다 많은 레지스터가 필요하면 초과분이 느린 로컬 메모리로 스필됩니다. 스필은 낮은 점유율보다 더 큰 성능 저하를 일으킬 수 있습니다.

레지스터나 공유 메모리를 줄이면 점유율이 높아지지만, 지나치게 줄이면 스필이 발생합니다. 최적 지점은 균형을 이루는 곳입니다.

튜닝하기 전에 항상 컴파일러가 보고하는 실제 레지스터 및 공유 메모리 사용량을 확인하십시오. 추측하면 대개 잘못된 조정 항목을 선택하게 됩니다.

레지스터와 공유 메모리는 SM의 고정된 예산이며, 가장 엄격한 제한이 점유율의 상한을 결정한다는 것을 확인했습니다. 스필에 주의하십시오. 🧮