Warp Divergence QUIZ - 추가로 조건이 스레드 단위가 아닌 블록 단위면 괜찮은 이유에 대해

🧩 Q1. 개념 체크

1.1. Warp divergence가 발생하는 상황을 한 줄로 설명해봐.
(“warp 내부 스레드들이 ~~~ 해서 GPU가 ~~~ 하게 되는 것” 이런 식으로)

1.2. divergence가 생기면 실행 시간 관점에서 어떤 일이 벌어지는지 설명해봐.

 

 

🧩 Q2. 코드 패턴 – divergence 여부 판단

다음 커널 조각에서 warp divergence가 생기는지/안 생기는지 각각 답해봐.

int lane = threadIdx.x & 31;
if (lane < 16) {
    x = x + 1.0f;
}
/////////////////////////////

int lane = threadIdx.x & 31;
if (blockIdx.x % 2 == 0) {
    x = x + 1.0f;
}
/////////////////////////////

if (input[gid] > 0.0f) {
    x = x * 2.0f;
} else {
    x = x - 2.0f;
}

각각에 대해:

• divergence 있음 / 없음
• 있다면 왜 있는지, 없으면 왜 없는지 짧게 써봐.

 

🧩 Q3. “숨겨지는” divergence

다음 두 상황 중, 어느 쪽이 실행 시간에서 divergence 영향이 더 잘 드러날지 생각해봐.

• (1) 전체 스레드가 많고, SM마다 warp가 수백 개씩 동시에 스케줄되는 큰 커널
• (2) 단 한 warp만 사용하는 아주 작은 커널 (grid=1, block=32)

어떤 경우에 divergence 비용이 “숫자로 눈에 더 잘 보이는지”와 이유를 같이 써봐.

 

 

🧩 Q4. branch-free로 바꾸기

다음 ReLU 코드가 있다고 하자:

if (x > 0.0f) {
    y = x;
} else {
    y = 0.0f;
}

4.1. 이 코드는 warp divergence를 유발할 가능성이 있을까? 있다면 어떤 상황에서일까?
4.2. 이걸 **branch-free(분기 없는 형태)**로 바꾸는 식을 떠올려봐. (힌트: mask, max, 곱셈 등)

 

 

🧩 Q5. “좋은 분기” vs “나쁜 분기”

다음 두 표현 중, 어떤 조건이 warp divergence 측면에서 더 안전한지 말해봐.

• (A) if (blockIdx.x % 2 == 0) {...}
• (B) if (threadIdx.x % 2 == 0) {...}

왜 그런지도 같이 설명해봐.
(힌트: warp 내 스레드들이 같은 값을 보냐, 다른 값을 보냐)

 

 

warp 단위 실행 모델에서의 조건 관측

warp 내부 스레드들이 모두 같은 분기 경로를 선택하면 warp divergence 가 없다.

block 단위 조건은 block 내부 모든 스레드에게 값은 값을 주기 때문에, warp 안 스레드 모두 같은 분기 선택

 

GPU 의 실행 단위 = 스레드가 아닌 warp

GPU 의 실제 실행 단위는 warp로 이는 32 개의 thread 묶음

warp 는 동일한 Progrram counter 를 공유한다.

warp 는 한 사이클에 모든 thread 가 같은 instruction 을 수행해야 한다.

 

warp 안에서 조건이 thread 마다 다르면 divergence 발생? Q1. 개념 체크

1.1. Warp divergence가 발생하는 상황을 한 줄로 설명해봐.

(“warp 내부 스레드들이 ~~~ 해서 GPU가 ~~~ 하게 되는 것” 이런 식으로)

 

1.2. divergence가 생기면 실행 시간 관점에서 어떤 일이 벌어지는지 설명해봐.

 

 

 

 

 

🧩 Q2. 코드 패턴 – divergence 여부 판단

다음 커널 조각에서 warp divergence가 생기는지/안 생기는지 각각 답해봐.

 

int lane = threadIdx.x & 31;

if (lane < 16) {

    x = x + 1.0f;

}

/////////////////////////////

 

int lane = threadIdx.x & 31;

if (blockIdx.x % 2 == 0) {

    x = x + 1.0f;

}

/////////////////////////////

 

if (input[gid] > 0.0f) {

    x = x * 2.0f;

} else {

    x = x - 2.0f;

}

각각에 대해:

 

divergence 있음 / 없음

있다면 왜 있는지, 없으면 왜 없는지 짧게 써봐.

 

 

🧩 Q3. “숨겨지는” divergence

다음 두 상황 중, 어느 쪽이 실행 시간에서 divergence 영향이 더 잘 드러날지 생각해봐.

 

(1) 전체 스레드가 많고, SM마다 warp가 수백 개씩 동시에 스케줄되는 큰 커널

(2) 단 한 warp만 사용하는 아주 작은 커널 (grid=1, block=32)

어떤 경우에 divergence 비용이 “숫자로 눈에 더 잘 보이는지”와 이유를 같이 써봐.

 

 

 

 

 

🧩 Q4. branch-free로 바꾸기

다음 ReLU 코드가 있다고 하자:

 

if (x > 0.0f) {

    y = x;

} else {

    y = 0.0f;

}

4.1. 이 코드는 warp divergence를 유발할 가능성이 있을까? 있다면 어떤 상황에서일까?

4.2. 이걸 **branch-free(분기 없는 형태)**로 바꾸는 식을 떠올려봐. (힌트: mask, max, 곱셈 등)

 

 

 

 

 

🧩 Q5. “좋은 분기” vs “나쁜 분기”

다음 두 표현 중, 어떤 조건이 warp divergence 측면에서 더 안전한지 말해봐.

 

(A) if (blockIdx.x % 2 == 0) {...}

(B) if (threadIdx.x % 2 == 0) {...}

왜 그런지도 같이 설명해봐.

(힌트: warp 내 스레드들이 같은 값을 보냐, 다른 값을 보냐)

 

 

 

 

 

warp 단위 실행 모델에서의 조건 관측

warp 내부 스레드들이 모두 같은 분기 경로를 선택하면 warp divergence 가 없다.

 

block 단위 조건은 block 내부 모든 스레드에게 값은 값을 주기 때문에, warp 안 스레드 모두 같은 분기 선택

 

 

 

GPU 의 실행 단위 = 스레드가 아닌 warp

 

GPU 의 실제 실행 단위는 warp로 이는 32 개의 thread 묶음

 

warp 는 동일한 Progrram counter 를 공유한다.

 

warp 는 한 사이클에 모든 thread 가 같은 instruction 을 수행해야 한다.

 

 

 

warp 안에서 조건이 thread 마다 다르면 divergence 발생, warp 안에서 조건이 동일하면 노상관

 

Warp = {t0,t1,t2,t3, ... t31}

조건:
t0: true
t1: false
t2: true
t3: false
...
=> warp 내부에서 섞임 => divergence 발생

////////////////////////////////////////
Block 0:
  Warp0: 모두 true
  Warp1: 모두 true
  Warp2: 모두 true

Block 1:
  Warp0: 모두 false
  Warp1: 모두 false
  Warp2: 모두 false