GPU memory transaction, Byte 단위 사고 정리

• 왜 GPU 커널에서 16 / 32 / 128 B 같은 숫자가 반복되는지이해
• lane - warp - memory transaction 으로 이어지는 실제 하드웨어 흐름 정리
• lane specialization = 메모리 transaction 설계 라는 말의 정확한 의미 고정

 

b = bit, B = byte( 8 bit )

 

GPU 는 byte 덩어리로만 메모리를 움직인다

 

GPU 메모리 시스템의 핵심 원칙

GPU 는 필요한 원소만큼 읽지 않는다 정해진 크기의 byte 덩어리 (transaction) 를 읽고, 그 안에서 필요한 데이터를 사용한다

• 연산 단위 - thread / warp
• 메모리 단위 - byte transaction

 

메모리 계층별 실제 단위

Global - L2 / L1

• transaction 크기 : 보통 128 B (warp 기준 최대 효율)
• 주소가 흩어지면 128 transaction 이 여러 개 발생

cp.async

• lane 단위 복사 크기 = 16 B
• 1 lane = 16 B
• 단일 warp 전체 =  최대 32 x 16B = 512B/

단, 512B 가 그대로 움직이는 것이 아닌, 내부에서 128 B transaction 여러 개로 분해됨

Shared memory

• bank 기반
• Ampere : 32 banks x 4B
• 여기서는 transaction 보다 bank conflict 가 더 중요

 

lane 과 warp 의 역할 분리

• lane = 주소생성기
• warp = 주소들을 모아 transaction 생성

즉

1. 각 lane 이 이 주소 읽고 싶다를 제시
2. 하드웨어가 warp 의 주소들을 분석
3. 가능한 한 연속된 byte 범위로 묶어서
4. 128 B / 32 B transaction 을 발사

 

warp coalescing 예시

lane 0 : addr 0
lane 1 : addr 4
...
lane 31: addr 124

• 주소 범위 : 0 ~ 127
• 128 transaction 1 개

낭비 발생

lane 0~15 : addr 0~60
lane 16~31: unused

• 실제 필요 데이터 64B
• 하지만 포함된 최소 cache line : 128B
• 128 transaction 1개 ( 64 B 낭비 )

GPU 는 이걸 신경쓰지 않음, 커널 작성자가 신경 써야 한다.

 

lane 에 16B 를 매핑한다의 정확한 의미

lane 이 16B 정렬된 주소를 담당하게 설계해서 warp 전체 요청이 128B 경계에 깔끔히 떨어지도록 만든다

cp.async 기준

• 명령어 자체가 16B 고정
• 따라서
  • lane 1개 = 16B

 

실제 커널 예시 해석

K tile 로드

• Kdim = 16
• dtype = half = 2B

16 x 2B = 32B (1 row)

• lane 2개가 한 row 담당

결과

• 16 rows x 32B = 512B
• warp 전체 주소가 연속
• 128B transaction x 4

 

transaction = GPU 메모리 시스템의 기본 동작 단위 - 연속, 정렬, 낭비없는 상태가 최적의 상태