coalescing / stride 접근 / warp 메모리 패턴 QUIZ

🧩 퀴즈 1 — Warp 메모리 접근 패턴

warp(32 threads)가 다음 주소를 읽는 상황:

 

thread 0 → addr A + 0
thread 1 → addr A + 4
thread 2 → addr A + 8
...
thread 31 → addr A + 124

❓ Q1. 이 접근은 coalesced인가, non-coalesced인가?

 

Coalesced,

float = 4bytes, 

warp 32 threads - 32 * 4 = 128 bytes

정확히 연속된 128 bytes, 딱 한 개의 DRAM 트랜잭션으로 읽힌다

 

 

🧩 퀴즈 2 — stride 접근 패턴

warp(32 threads)가 다음처럼 stride = 8 로 읽는다고 하자:

thread i → addr base + (i * 8 * 4 bytes)

여기서 4bytes는 float 크기.

❓ Q2. thread 0~31 이 실제 읽는 바이트 오프셋 범위를 써라

(즉, threadIdx.x * (stride * sizeof(float)))

 

32 * 31

❓ Q2-2. 이 패턴은 coalesced될 수 있는가?

(정답: Yes / No)

 

no, GPU warp 메모리 coalescing 은 하나의 128B 라인 내에서 thread 32 개가 묶여 읽을 때 발생, 

stride = 8, 모든 thread 주소가 128B line 을 다 뚫고 퍼짐

warp 가 7~8 개 이상의 서로 다른 DRAM line 에 걸쳐 읽게 됨

 

 

🧩 퀴즈 3 — DRAM 트랜잭션 개수 예측 (개념적)

warp 1개가 다음 두 형태로 global memory를 읽는다.

1. 연속(read[i])
2. stride=16(read[i*16])

GPU는 global memory를 보통 128B 단위로 트랜잭션한다고 가정하자.

warp(32 threads)가 float(4B) 읽을 때:

❓ Q3-1. coalesced (연속) 접근에서는 몇 개의 128B 트랜잭션이 필요한가?

 

32 * 4 의 정확히 1 트랜잭션

 

❓ Q3-2. stride=16 일 때, warp가 총 몇 개의 트랜잭션을 만들 가능성이 높은가?

(정확한 숫자 말고 “대략 몇 배 이상 더 많아진다” 정도로 답해도 됨)

 

 

🧩 퀴즈 4 — stride 큰 경우 튀어나온 결과 해석

너가 실험에서 나온 결과:

coalesced: ~0.33 ms
stride=8 ~1.49 ms
stride=16 ~1.47 ms
stride=32 ~1.47 ms

❓ Q4. stride=8 이후로 stride=16, stride=32 가 더 나빠지지 않는 이유를 설명해라.

(힌트: warp가 이미 “각 스레드가 모두 다른 128B 라인”을 읽게 되면…)

 

충분히 나빠짐 각 thread 가 다른 DRAM line 을 읽는 상태

 

🧩 퀴즈 5 — 주소 패턴 비교

다음 두 접근 패턴이 있다:

thread 0 → i
thread 1 → i+1
thread 2 → i+2
...
thread 31 → i+31

////////////////
thread 0 → i
thread 1 → i+64
thread 2 → i+128
...
thread 31 → i+1984

❓ Q5-1. 어떤 패턴이 더 빠르고, 왜 그런가?

❓ Q5-2. B 패턴을 빠르게 만들려면 어떤 구조적 변화(데이터 layout)를 해야 할까?