Non-Uniform L2 Cache Latency Across the Streaming Multiprocessors of an NVIDIA L40

기술 요약: NVIDIA L40의 비균일 L2 캐시 지연 시간

문제 정의
전체 AD102 구현을 기반으로 하는 NVIDIA L40 GPU는 관습적으로 성능 분석 도구 및 학술 문헌에서 단일한 디바이스 전역 히트 지연 시간을 가진 균일한 풀(pool)로 모델링되는 96 MiB L2 캐시를 노출한다. 이러한 추상화는 L2 캐시에 접근하는 모든 스트리밍 멀티프로세서(SM)가 동일한 비용을 발생시킨다는 가정을 전제로 한다. 그러나 물리적 실체로서의 L2 캐시는 메모리 컨트롤러에 부착된 여러 개의 슬라이스(slice)로 구성되며, 온칩 네트워크(NoC)를 통해 접근된다. V100, A100, H100 아키텍처에 대한 선행 연구들은 NoC 지연 시간이 비균일하다는 점을 확립했으나, 해당 연구들은 SM당 평균 지연 시간은 대략 균일하며 비균일성은 주로 파티션 경계에서의 불연속적인 단계(discrete steps)로 나타난다는 점을 강조했다. 단일 다이(die), 단일 파티션 구조인 Ada Lovelace 데이터 센터 GPU(L40과 같은)가 이와 유사한 구조를 보이는지, 그리고 만약 그렇다면 SM당 평균 지연 시간이 균일한지 아니면 배치에 따라 달라지는지는 여전히 미해결 과제로 남아 있었다.

방법론
저자들은 높은 권한이나 하드웨어 성능 카운터 없이도 모든 물리적 SM으로부터 L2 히트 지연 시간을 측정할 수 있는 사용자 수준의 마이크로벤치마킹 방법론을 개발했다.

• 프로브 설계: 측정에는 단일 스레드, 단일 워프 종속 로드 체인(ld.global.cg)을 사용하여 캐시 패브릭을 통한 단일 왕복(round-trip)을 보장한다. 두 가지 버퍼 형태가 사용되었다: 특정 128 B 슬라이스를 조사하기 위한 "단일 라인(single-line)" 형태와 모든 슬라이스로의 평균 거리를 측정하기 위한 "체인(chain)" 형태(2 MiB 무작위 순열)이다.
• SM 격리: 온칩 경합을 제거하기 위해 전역 턴 카운터(global turn counter)를 사용하여 블록들을 직렬화한다. 각 블록은 atomicAdd를 통해 자신의 차례를 점유하고, 대기한 후, 웜업(warm-up)을 수행하고, 격리된 상태에서 종속 로드를 타이밍 측정한다.
• 커버리지: 프로브는 L40의 142개 물리적 SM과 256개 슬라이스 프로브를 모두 커버한다. 일반화 테스트를 위해 유사한 프로브가 NVIDIA RTX 5090(Blackwell GB202)에서도 실행되었다.
• 모델링: 지연 시간 데이터는 다음과 같은 가산적 비균일 캐시 접근(NUCA) 모델을 사용하여 분석된다: $L = \mu + a(sm) + b(slice)$, 여기서 $a(sm)$은 SM 배치 항을, $b(slice)$는 슬라이스 배치 항을 나타낸다.

주요 결과

1. 유의미한 지연 시간 변동: 균일 모델과 달리, L40의 L2 히트 지연 시간은 일정하지 않다. 이는 222.5에서 339.2 사이클에 걸쳐 분포하며(52% 범위), SM 간 평균 변동은 57.2 사이클이다. 노이즈 플로어는 무시할 수 있는 수준(<0.01 사이클)으로, 이 변동이 물리적 현상임을 확인시켜 준다.
2. 이중 대칭성: SM 배치 항은 강력한 이중 대칭성을 보인다. 식별자 72를 기준으로 142개의 SM을 나누었을 때, 거의 동일한 지연 시간 프로파일(상관관계 $r=0.999$)을 가진 두 그룹(72개 및 70개 SM)으로 나뉜다. 이는 두 개의 대칭적인 6-GPC 클러스터로 구성된 AD102 플로어플랜과 일치한다.
3. 연속적 구배 vs 불연속적 단계: 파티션 경계에서 불연속적인 지연 시간 단계를 보여주는 A100/H100과 달리, L40은 이봉형 간극(bimodal gap) 없이 SM 전체에 걸쳐 연속적인 지연 시간 구배(gradient)를 나타낸다.
4. 가산 모델 적합도: 가산 모델은 지연 시간 변동의 87%를 설명한다($R^2=0.87$). 단일 랭크-1 상호작용 항을 추가하면 이 수치가 0.98까지 증가하며, 이는 지연 시간 맵이 GPC 레이아웃을 따르는 저차원 기하학적 구조임을 나타낸다.
5. 일반화 및 특이성:
   • 교차 아키텍처: RTX 5090 역시 비균일성(46% 확산)을 보이지만, 대칭성이 더 약하며($r=0.80$) 절대적인 지연 시간은 더 느리다.
   • 디바이스 핑거프린팅: 물리적으로 구별되는 동일 모델의 두 L40 장치는 거의 동일한 평균 지연 시간(0.28 사이클 차이)에도 불구하고, 각 SM별 맵을 기반으로 100% 정확도로 구별 가능하다. 즉, per-SM 맵은 안정적인 하드웨어 정체성 역할을 한다.
6. 안정성: per-SM 지연 시간 맵은 두 장치 모두에서 1시간 동안 100% 사용률을 유지한 후에도 변하지 않았으며, 이는 토폴로지가 일시적인 상태가 아니라 실리콘의 고유한 속성임을 확인시켜 준다.

기여 및 의의
본 논문은 다음과 같은 기여를 주장한다:

• 토폴로지 매핑: L40에 대한 최초의 per-SM L2 지연 시간 맵을 제공하여, 균일 캐시 가정이 지연 시간에 민감한 워크로드에 대해 유효하지 않음을 입증(52%의 확산 확인)한다.
• 자기 위치 식별 프리미티브(Self-Localization Primitive): 저자들은 단일 L2 지연 시간 프로브를 통해 커널이 자신의 물리적 SM 배치를 99.2% 정확도로(혼합 장치 세트에서는 92%) 식별할 수 있으며, 특정 디바이스 다이를 100% 정확도로 식별할 수 있음을 보여준다. 이는 피해 커널으로부터 데이터를 추출하는 것이 아니라 커널 자신의 배치를 추출하는 것이므로, 공격이 아닌 "자기 위치 식별" 프리미티브로 프레임화된다.
• 성능 영향: 지연 시간 맵에 따라 작업을 분배함으로써(더 빠른 SM에 작업 배치), 저자들은 L2 상주 워크로드에서 메이크스팬(makespan)을 10.9% 감소시켰다. 이 이점은 DRAM 중심 영역에서는 사라지며, 이는 토폴로지의 영향이 레지임(regime)에 따라 특이적임을 보여준다.
• 레지던시(Residency)의 재해 해석: 기존의 단일 스레드 레지던시 제어(용량 경계, 라인 태그 스케일링, 지속되는 L2 결과 등)를 상수로서가 아니라, 이 비균일 분포에 대한 평균값으로 재해석한다.

범위 및 한계
저자들은 보안 관련 주장의 범위를 명시적으로 제한한다. 이 기술은 자기 위치 식별(커널이 자신의 배치를 읽는 것) 및 디바이스 핑거프린팅을 가능하게 한다. 이는 별도의 공존하는 "피해" 커널의 배치를 추론하거나 비밀 데이터를 추출하는 것이 아니다. 그러한 기능은 본 연구에서 측정되지 않은 경합 또는 간섭 채널을 필요로 한다. 본 연구는 두 대의 L40과 한 대의 RTX 5090으로 제한된다. 또한 %smid와 물리적 플로어플랜 간의 매핑은 다이 사진을 통해 독립적으로 확인된 것이 아니라 대칭성과 GPC 수를 통해 추론된 것이다. 절대적인 사이클 수는 마이크로벤치마크 설계에 특화된 것이며, 독립적인 요청과 L1 사용이 발생하는 실제 애플리케이션 처리량으로 직접 변환되지는 않는다.