이 글에서는 AMD의 RDNA 4 GPU 아키텍처에 도입된 혁신적인 동적 레지스터 할당 기능에 대해 논의합니다. 이 기능은 점유율과 레지스터 수 간의 균형을 맞추는 데 중점을 두고 있습니다. RDNA 4의 명령어 집합 아키텍처(ISA)는 최대 256개의 벡터 범용 레지스터(VGPR)에 접근할 수 있도록 하며, 각 레지스터는 wave32 모드에서 1024비트 폭을 가집니다. 이 아키텍처는 SIMD당 192KB의 레지스터 파일을 특징으로 하여, 96개 이하의 벡터 레지스터로 최대 점유율을 가능하게 합니다.
동적 VGPR 할당은 스레드가 최소 VGPR 할당으로 시작하고 실행 중에 이를 조정할 수 있도록 하여 스레드 수준의 병렬성을 향상시키고 레이 트레이싱 작업에서 지연 민감도를 줄입니다. 이 메커니즘은 활성 스레드 수를 설정하고 동적 VGPR 매개변수를 제어하는 칩 전체의 `SQ_DYN_VGPR` 레지스터를 포함합니다. 그러나 할당 요청이 실패할 수 있으며, 이로 인해 바쁜 대기 상태가 발생할 수 있어 성능에 영향을 미칠 수 있습니다. AMD는 여러 스레드의 동시 레지스터 할당 요청으로 인해 발생할 수 있는 교착 상태를 완화하기 위해 교착 상태 회피 모드를 구현했습니다.
현재 동적 VGPR 모드는 wave32 컴퓨트 셰이더에만 제한되며, 그래픽 셰이더나 wave64 셰이더와 함께 사용할 수 없습니다. 비교적으로, NVIDIA의 동적 레지스터 할당 메커니즘인 `setmaxnreg`는 레지스터 재할당을 허용하지만 스레드 간의 동기화를 요구하여 유연성을 제한할 수 있습니다. Intel의 GPU는 전통적으로 고정 레지스터 할당을 사용하여 동적 할당이 그들의 아키텍처에 덜 유익합니다.
전반적으로 AMD의 동적 VGPR 할당은 레이 트레이싱 및 일반 컴퓨트 작업의 효율성을 향상시키는 중요한 발전으로, 레지스터 파일 용량을 늘리지 않고도 더 많은 스레드가 활성화될 수 있도록 합니다. 이 기능은 AMD가 성능을 개선하고 GPU 아키텍처의 한계를 해결하기 위한 지속적인 노력을 반영합니다.
* 이 글은
old.chipsandcheese.com의 기사를 요약한 것입니다. 전체 기사의 내용은
이곳에서 확인하실 수 있습니다.