1. 연산 개요
Selection은 조건에 따라 두 개 이상의 후보 값 중 하나를 출력으로 선택하는 연산이다.
가장 기본적인 이항 Selection은 다음과 같다.
y[i] = predicate[i] ? a[i] : b[i]
이를 조건문 형태로 표현하면 다음과 같다.
if predicate[i] is true:
y[i] = a[i]
else:
y[i] = b[i]
데이터 의존성은 다음과 같다.
predicate[i] ─┐
a[i] ─────────┼→ Select → y[i]
b[i] ─────────┘
Selection의 출력은 다음 세 값에 의존한다.
1. 조건 predicate[i]
2. true일 때 선택할 값 a[i]
3. false일 때 선택할 값 b[i]
Selection은 조건에 따라 실행 경로 자체를 나누는 Branch와 다르다.
Selection은 일반적으로 두 후보 값을 데이터로 취급하고, predicate에 따라 최종 결과를 결정한다.
Branch:
어떤 명령 경로를 실행할지 결정
Selection:
이미 정의된 후보 값 중 어떤 값을 출력할지 결정
Selection은 다음 연산의 기초가 된다.
- ReLU
- Clamp
- Maximum
- Minimum
- Thresholding
- Mask 적용
- 조건부 값 대체
- Dropout 출력 선택
- Attention masking
- NaN 처리
- Boundary value 처리
- Quantization saturation
- Piecewise function
- 조건부 update
- Conditional move
2. 기본 수학적 정의
Predicate p, true 후보 a, false 후보 b가 있다고 하자.
Selection은 다음과 같이 정의된다.
select(p, a, b)
=
a, if p is true
b, if p is false
Elementwise 형태에서는 다음과 같다.
y[i] = select(p[i], a[i], b[i])
또는:
y[i] = p[i] ? a[i] : b[i]
각 출력 원소는 같은 위치의 predicate와 두 후보 값에 의존한다.
p[0], a[0], b[0] → y[0]
p[1], a[1], b[1] → y[1]
p[2], a[2], b[2] → y[2]
다른 위치의 결과에는 의존하지 않는다.
따라서 Elementwise Selection은 원소 단위 병렬 실행이 가능하다.
3. Selection의 논리적 의미
Selection은 predicate가 표현하는 조건의 의미를 데이터 출력으로 반영한다.
예를 들어:
p[i] = x[i] > 0
y[i] = select(p[i], x[i], 0)
은 다음과 같다.
x[i] > 0이면 x[i]
그렇지 않으면 0
즉 ReLU다.
y[i] = max(x[i], 0)
다른 예:
p[i] = x[i] < lower
y[i] = select(p[i], lower, x[i])
이는 lower bound를 적용하는 연산이다.
Selection 자체는 후보 값의 의미를 알지 못한다.
Selection이 아는 것:
predicate가 true인지 false인지
Selection이 알지 못하는 것:
왜 a와 b 중 하나를 선택하는지
고수준 의미는 predicate와 후보 값의 구성으로 결정된다.
4. Scalar Selection
후보 중 하나 또는 둘 모두가 scalar일 수 있다.
True 후보가 scalar인 경우
y[i] = p[i] ? c : x[i]
False 후보가 scalar인 경우
y[i] = p[i] ? x[i] : c
두 후보 모두 scalar인 경우
y[i] = p[i] ? high : low
예를 들어 boolean mask를 0과 1의 수치 mask로 변환할 수 있다.
m[i] = p[i] ? 1 : 0
이 경우 predicate의 논리적 값이 일반 register의 수치값으로 materialize된다.
5. Tensor-Tensor Selection
두 후보가 모두 tensor인 경우다.
y[i] = p[i] ? a[i] : b[i]
예:
P = [true, false, true]
A = [10, 20, 30]
B = [1, 2, 3]
결과:
Y = [10, 2, 30]
기본 실행 구조는 다음과 같다.
Index
→ Load Predicate
→ Load A
→ Load B
→ Select
→ Store Y
하지만 predicate가 같은 kernel 안에서 생성되면 별도의 predicate load가 필요하지 않을 수 있다.
Load Comparison Inputs
→ Compare
→ Predicate Register
→ Load/Compute A
→ Load/Compute B
→ Select
→ Store Y
6. Broadcasting Selection
Predicate 또는 후보 값에 broadcasting이 적용될 수 있다.
예를 들어 channel별 조건을 전체 batch에 적용할 수 있다.
y[n, c] =
predicate[c] ? a[n, c] : b[n, c]
또는 scalar 조건 하나가 전체 tensor를 선택할 수 있다.
Y = global_predicate ? A : B
Broadcasting이 포함되면 다음 관계를 기록해야 한다.
어떤 축으로 predicate가 반복되는가
어떤 후보가 scalar인가
어떤 후보가 tensor인가
실행 관점에서는 broadcast 값의 재사용과 index 계산이 중요해진다.
7. 입력과 출력 domain
Selection의 입력은 일반적으로 다음으로 구성된다.
P: predicate 또는 boolean domain
A: value domain
B: value domain
출력은 후보와 호환되는 domain을 가진다.
Y domain = common_domain(A, B)
가장 단순한 경우:
dtype(A) = dtype(B) = dtype(Y)
하지만 후보의 dtype이 다르면 type promotion이나 conversion이 필요할 수 있다.
예:
A: FP32
B: INT32
Y: FP32
실제 연산은 다음을 포함할 수 있다.
1. B를 FP32로 변환
2. Predicate에 따라 A 또는 변환된 B 선택
따라서 Selection의 의미 명세에는 후보의 dtype과 변환 규칙을 포함해야 한다.
8. 데이터 의존성
기본 Selection은 다음과 같다.
y[i] = select(p[i], a[i], b[i])
출력은 세 입력에 논리적으로 의존한다.
p[i]는 어떤 후보를 사용할지 결정한다.
a[i]는 true일 때 출력이 된다.
b[i]는 false일 때 출력이 된다.
그러나 실제 실행에서 두 후보가 모두 계산되어야 하는지는 별개의 문제다.
후보가 이미 존재하는 값인 경우
a[i], b[i]가 이미 register 또는 memory에 존재
Selection은 단순히 둘 중 하나를 선택한다.
후보가 계산식인 경우
y[i] = p[i] ? f(x[i]) : g(x[i])
다음 두 실행 방식이 가능하다.
방식 1:
f와 g를 모두 계산한 뒤 Select
방식 2:
Predicate에 따라 한쪽 경로만 실행
따라서 Selection의 수학적 표현과 실제 계산 비용은 같지 않을 수 있다.
9. 의미 불변성
Selection 구현이 달라져도 다음 조건은 유지되어야 한다.
9.1 Predicate 의미
Predicate가 true일 때 true 후보가 선택되어야 한다.
p = true → y = a
p = false → y = b
후보 순서를 바꾸려면 predicate를 반전해야 한다.
select(p, a, b)
=
select(not p, b, a)
9.2 후보 대응 관계
Elementwise Selection에서는 같은 논리적 위치의 후보를 선택해야 한다.
y[i] = select(p[i], a[i], b[i])
다른 위치의 값을 사용하면 다른 연산이 된다.
9.3 출력 cardinality
각 논리적 위치마다 출력 하나가 생성된다.
one predicate position
+
two candidate positions
→
one output position
9.4 비선택 후보의 의미
수학적으로는 선택되지 않은 후보가 출력에 영향을 주지 않는다.
p = true인 경우 b는 출력값에 영향을 주지 않는다.
그러나 실제 프로그램에서는 비선택 후보를 계산하는 과정이 side effect를 가질 수 있다.
따라서 순수 함수와 side effect를 가진 계산을 구분해야 한다.
10. Selection과 Branch의 차이
Selection과 Branch는 모두 조건을 사용하지만 역할이 다르다.
Selection
y = p ? a : b
핵심은 결과값 선택이다.
두 후보 값
→ 하나의 출력값
Branch
if p:
execute path A
else:
execute path B
핵심은 실행할 명령 경로 선택이다.
두 명령 경로
→ 하나 또는 여러 실행 효과
Selection은 일반적으로 side effect가 없는 값 계산과 잘 맞는다.
Branch는 다음처럼 복잡한 경로를 포함할 수 있다.
- 여러 산술 연산
- Memory store
- Loop
- Function call
- Atomic update
- Synchronization
따라서 모든 Branch를 Selection으로 바꿀 수 있는 것은 아니다.
11. Eager Selection과 Lazy Branch
다음 표현을 생각하자.
y = p ? f(x) : g(x)
Eager Selection
두 후보를 모두 계산한다.
a = f(x)
b = g(x)
y = select(p, a, b)
비용:
cost(f) + cost(g) + cost(select)
Lazy Branch
선택된 경로만 계산한다.
if p:
y = f(x)
else:
y = g(x)
이상적인 비용:
p가 true이면 cost(f)
p가 false이면 cost(g)
GPU에서는 warp divergence 때문에 warp 내 일부 thread가 양쪽 경로를 모두 실행할 수 있다.
따라서 실제 비용은 predicate 분포와 경로 복잡도에 따라 달라진다.
12. Branchless Selection
짧은 조건부 연산은 branch 없이 selection instruction으로 표현할 수 있다.
p = x > 0
y = select(p, x, 0)
추상 구조:
Compare
→ Predicate
→ Select
장점:
- 명시적 branch 제거
- Warp reconvergence 부담 감소
- 짧은 조건식에서 예측 가능한 실행
- Predicate를 register에 유지 가능
단점:
- 두 후보 계산이 모두 필요할 수 있음
- 비선택 후보의 memory load가 발생할 수 있음
- 복잡한 후보식에서는 불필요한 연산 증가
Branchless는 무조건적인 성능 향상이 아니라 하나의 실행 선택이다.
13. Arithmetic Mask 방식
Selection을 0과 1의 mask를 이용한 산술식으로 표현할 수 있다.
m ∈ {0, 1}
y = m*a + (1-m)*b
m = 1이면:
y = a
m = 0이면:
y = b
하지만 이 방식은 일반 Selection과 완전히 같은 의미를 가지지 않을 수 있다.
추가 연산
Multiply 두 번
Add 한 번
Mask 변환
특수값 문제
예를 들어 m = 0, a = Inf이면:
0 * Inf = NaN
따라서:
m*a + (1-m)*b
는 원래 Selection이 b를 반환해야 하는 상황에서도 NaN을 만들 수 있다.
즉 다음 두 식은 일반적인 유한값에서는 같을 수 있지만 IEEE 특수값에서는 다르다.
select(m, a, b)
m*a + (1-m)*b
14. Bitwise Mask Selection
정수나 bit pattern에서는 mask를 이용한 bitwise selection이 가능하다.
mask = all_ones, if p is true
mask = all_zero, if p is false
Selection은 다음처럼 표현할 수 있다.
y = (mask & a) | (~mask & b)
장점:
- Branch 없이 구현 가능
- Integer와 bit pattern 선택에 적합
- SIMD/SIMT mask와 연결 가능
제한:
- Floating-point 값을 사용할 경우 bit representation 단위로 처리해야 함
- Mask representation이 정확해야 함
- 추가 bitwise instruction이 필요할 수 있음
- NaN payload나 signed zero를 bitwise하게 그대로 보존하지만 수치 의미와는 별개
15. Predicate Register Selection
GPU에서는 comparison 결과가 predicate register에 저장되고, Selection이 이를 직접 사용할 수 있다.
개념적으로:
P0 = x > 0
Rout = select(P0, Rx, Rzero)
이 구조에서는 predicate를 일반 register나 global memory로 변환할 필요가 없다.
Compare
→ Predicate
→ Select
이것은 논리적 중간값 materialization 제거의 대표적인 사례다.
16. Selection과 Materialization
다음과 같이 Comparison 결과를 먼저 mask tensor로 저장할 수 있다.
mask[i] = x[i] > 0
y[i] = mask[i] ? a[i] : b[i]
분리 구조:
Compare kernel
→ mask global store
→ mask global load
→ Select kernel
융합 구조:
Load X
→ Compare
→ Predicate register
→ Select
→ Store Y
Mask가 다른 곳에서 사용되지 않는다면 다음을 제거할 수 있다.
- Mask tensor allocation
- Mask global store
- Mask global load
- 추가 kernel launch
- 추가 index 계산
17. Selection과 관측 가능성
중간 predicate가 외부에서 관측되는지 여부가 fusion 가능성을 결정한다.
비관측 predicate
predicate가 Selection 하나에서만 사용됨
이 경우 register predicate로 유지할 수 있다.
관측 predicate
predicate 자체가 사용자 출력임
여러 consumer가 사용함
backward에서 필요함
debugging에 사용됨
이 경우 materialization이 필요할 수 있다.
Selection은 값의 관측 가능성뿐 아니라 predicate의 관측 가능성도 분석해야 한다.
18. Selection의 대수적 성질
18.1 동일 후보 제거
select(p, a, a) = a
Predicate와 Selection을 모두 제거할 수 있다.
단, predicate 계산이 side effect를 가지지 않는 순수 계산이어야 한다.
18.2 상수 Predicate
select(true, a, b) = a
select(false, a, b) = b
Compiler가 predicate를 compile time에 알면 불필요한 후보와 Selection을 제거할 수 있다.
18.3 Predicate 반전
select(p, a, b)
=
select(not p, b, a)
이 성질은 operand 순서와 조건 관계를 조정하는 데 사용된다.
18.4 Nested Selection
다음 중첩 구조를 생각할 수 있다.
select(p, select(q, a, b), c)
조건 영역을 분석하면 piecewise function으로 해석할 수 있다.
if p and q:
a
else if p:
b
else:
c
Predicate 관계에 따라 중복 비교나 후보를 제거할 수 있다.
19. Selection 분배와 재배치
순수 함수 f가 있을 때 다음을 비교할 수 있다.
f(select(p, a, b))
select(p, f(a), f(b))
수학적으로 f가 deterministic하고 side effect가 없다면 결과는 같을 수 있다.
하지만 실행 비용은 다르다.
첫 번째:
Select 한 번
→ f 한 번
두 번째:
f 두 번
→ Select 한 번
따라서 일반적으로 공통 후속 연산을 Selection 밖으로 이동하면 계산량을 줄일 수 있다.
select(p, f(a), f(b))
→
f(select(p, a, b))
다만 다음 조건을 검토해야 한다.
- f가 두 후보에 동일하게 적용되는가
- Precision과 rounding 순서가 같은가
- f가 NaN, Inf를 다르게 처리하지 않는가
- Side effect가 없는가
20. 공통 연산 Hoisting
다음 구조를 생각하자.
if p:
y = f(x) + c
else:
y = g(x) + c
공통 Add를 밖으로 이동할 수 있다.
t = select(p, f(x), g(x))
y = t + c
이 변환은 중복 연산을 줄일 수 있다.
반대로:
if p:
y = a*x
else:
y = b*x
는 다음과 같이 바꿀 수 있다.
scale = select(p, a, b)
y = scale*x
즉 후보 계산의 공통 구조를 factorization할 수 있다.
하지만 부동소수점에서는 연산 순서와 rounding이 달라질 수 있으므로 bitwise equality는 별도로 확인해야 한다.
21. ReLU
ReLU는 Selection의 대표적인 사례다.
y[i] = max(0, x[i])
Comparison과 Selection으로 표현하면:
p[i] = x[i] > 0
y[i] = select(p[i], x[i], 0)
추상 구조:
Load X
→ Compare with 0
→ Select X or 0
→ Store Y
가능한 lowering은 다음과 같다.
Comparison + Branch
Comparison + Selection
Direct Maximum
일반 유한값에서는 같은 결과를 만들 수 있지만 다음에서 차이가 있을 수 있다.
- NaN
- Signed zero
- Fast-math
- Maximum instruction의 NaN 규칙
22. Clamp
Clamp는 두 번의 조건 선택으로 표현할 수 있다.
기본 정의:
y[i] = min(max(x[i], lower), upper)
Selection 형태:
t[i] =
select(x[i] < lower, lower, x[i])
y[i] =
select(t[i] > upper, upper, t[i])
또는 세 구간 piecewise function으로 표현할 수 있다.
if x[i] < lower:
y[i] = lower
else if x[i] > upper:
y[i] = upper
else:
y[i] = x[i]
가능한 lowering:
두 Comparison + 두 Selection
FMAX + FMIN
Branch chain
23. Maximum과 Minimum
Maximum은 Selection으로 표현할 수 있다.
max(x, z)
=
select(x >= z, x, z)
Minimum은:
min(x, z)
=
select(x <= z, x, z)
하지만 floating-point에서는 NaN 처리 규칙을 확인해야 한다.
예를 들어:
x = NaN
z = 1
일 때:
select(x >= z, x, z)
에서는 comparison이 false이므로 z가 선택될 수 있다.
반면 NaN 전파형 maximum은 NaN을 반환해야 할 수 있다.
따라서 direct max/min과 Comparison-Selection이 항상 같은 것은 아니다.
24. Thresholding
Thresholding은 predicate에 따라 상수나 입력값을 선택한다.
Binary threshold:
y[i] =
high, if x[i] >= threshold
low, otherwise
Selection 표현:
p[i] = x[i] >= threshold
y[i] = select(p[i], high, low)
Hard threshold:
y[i] =
x[i], if abs(x[i]) >= threshold
0, otherwise
Selection은 quantization, pruning, sparse activation 생성에 사용될 수 있다.
25. Masked Fill
Masked Fill은 mask가 true인 위치를 특정 값으로 대체한다.
y[i] =
fill_value, if mask[i] is true
x[i], otherwise
Selection 형태:
y[i] = select(mask[i], fill_value, x[i])
Attention mask에서는 다음이 사용될 수 있다.
fill_value = -Inf
masked_score[i, j]
=
select(blocked[i, j], -Inf, score[i, j])
이는 additive mask와 유사한 효과를 만들지만 특수값 처리와 실행 구조가 다르다.
26. Attention Masking
Causal attention에서는 위치 비교 결과로 score를 선택할 수 있다.
blocked[i, j] = j > i
masked_score[i, j]
=
select(blocked[i, j], -Inf, score[i, j])
추상 구조:
Index Compare
→ Predicate
→ Select -Inf or Score
→ Softmax
이 경우 mask tensor를 별도로 만들 필요 없이 index comparison과 Selection을 attention score kernel 안에 fusion할 수 있다.
score 계산
→ causal index comparison
→ Select score or -Inf
→ 후속 Softmax 상태 계산
27. Dropout
Dropout은 mask에 따라 입력을 유지하거나 제거한다.
y[i] =
x[i] * scale, if keep[i] is true
0, otherwise
Selection 형태:
y[i] =
select(keep[i], x[i]*scale, 0)
Multiply mask 방식:
y[i] = x[i] * mask[i] * scale
일반적인 유한 입력에서는 비슷할 수 있지만 x = Inf이고 mask가 0이면 차이가 발생한다.
Selection:
select(false, Inf*scale, 0)
구현이 true 후보를 실제로 계산하지 않거나 값 선택 semantics를 유지하면 0을 반환할 수 있다.
Multiply mask:
Inf * 0 = NaN
따라서 Dropout의 정확한 특수값 의미에 따라 Select와 Multiply 방식이 다를 수 있다.
28. Conditional Update
기존 값을 조건부로 갱신할 수 있다.
new_value[i] =
candidate[i], if update[i] is true
old_value[i], otherwise
Selection 형태:
new_value[i]
=
select(update[i], candidate[i], old_value[i])
이 구조는 다음에서 사용된다.
- Optimizer parameter update
- Recurrent state update
- Early-exit state 유지
- Iterative algorithm
- Sparse update
- Masked assignment
In-place 형태:
state[i] =
select(update[i], candidate[i], state[i])
기존 state[i]가 false 후보이자 출력 buffer가 된다.
29. Piecewise Function
Selection은 여러 구간으로 정의된 함수를 표현한다.
예를 들어 Leaky ReLU:
y[i] =
x[i], if x[i] >= 0
alpha*x[i], otherwise
Selection 표현:
p[i] = x[i] >= 0
y[i] = select(p[i], x[i], alpha*x[i])
Hard Sigmoid와 Hard Swish 같은 함수도 여러 Selection과 affine transform의 조합으로 표현할 수 있다.
Comparison
→ Affine
→ Selection
→ Clamp
따라서 Selection은 piecewise function의 핵심 primitive다.
30. Selection과 Side Effect
수학적 Selection은 선택되지 않은 후보가 결과에 영향을 주지 않는다.
하지만 프로그램 표현에서는 후보 계산이 side effect를 가질 수 있다.
예:
y = p ? atomic_add(A) : atomic_add(B)
두 후보를 모두 계산한 뒤 Selection하면 원래 의미와 달라진다.
atomic_add(A)와 atomic_add(B)가 모두 실행됨
따라서 Eager Selection 변환은 다음 조건에서만 안전하다.
- 후보 계산이 순수함
- Memory store가 없음
- Atomic operation이 없음
- Exception이나 trap 의미가 없음
- Random state 변경이 없음
- Synchronization이 없음
31. 예외와 정의되지 않은 연산
선택되지 않은 후보가 정의되지 않은 값을 만들 수 있는 경우도 중요하다.
예:
y =
x != 0 ? 1/x : 0
Branch semantics에서는 x = 0일 때 division을 실행하지 않을 수 있다.
하지만 eager하게 두 후보를 모두 계산하면:
a = 1/x
b = 0
y = select(x != 0, a, b)
x = 0에서 division by zero가 발생한다.
부동소수점에서는 Inf가 생성될 수 있고, 다른 언어나 타입에서는 exception 또는 undefined behavior가 발생할 수 있다.
따라서 Branch-to-Selection 변환은 후보 계산의 정의 가능성까지 고려해야 한다.
32. Memory Load와 Selection
다음 표현을 생각하자.
y[i] =
p[i] ? A[index_a[i]] : B[index_b[i]]
Selection instruction 자체는 단순해도 두 후보를 모두 먼저 load하면 memory traffic이 증가한다.
Eager load:
load A[index_a]
load B[index_b]
select
Branch load:
if p:
load A[index_a]
else:
load B[index_b]
Predicate 분포와 memory locality에 따라 어느 방식이 유리한지 달라진다.
특히 후보 address가 유효하지 않을 가능성이 있다면 비선택 후보 load를 수행해서는 안 된다.
33. Safe Memory Selection
다음 코드는 위험할 수 있다.
value_a = A[index_a]
value_b = B[index_b]
y = select(p, value_a, value_b)
p = true여도 index_b가 유효하지 않으면 이미 invalid load가 발생한다.
따라서 address 자체를 선택한 뒤 한 번만 load하는 방식이 필요할 수 있다.
address =
select(p, address_a, address_b)
y = load(address)
이 경우 Selection이 값이 아니라 pointer 또는 address에 적용된다.
Select Address
→ Load Selected Address
이 구조는 memory safety와 load 수를 개선할 수 있지만 coalescing과 address divergence에 영향을 줄 수 있다.
34. Selection과 Warp Divergence
Selection 자체는 일반적으로 control-flow branch를 만들지 않으므로 warp divergence를 직접 발생시키지 않을 수 있다.
thread별 predicate가 다르더라도
모든 thread가 같은 Select instruction을 실행
그러나 다음 요소는 여전히 thread별로 다르다.
- 선택되는 값
- 후속 데이터
- 선택된 address
- 이후 계산 결과
Branch와 비교하면:
Selection:
같은 instruction stream
다른 데이터 결과
Branch:
서로 다른 instruction path
다만 compiler가 Selection을 Branch로 lowering하거나 후보 계산을 조건부 경로로 분리하면 divergence가 발생할 수 있다.
35. Selection의 추상 실행 모티프
Predicate가 이미 존재하는 경우
Load/Compute A
→ Load/Compute B
→ Select using Predicate
→ Store Y
Comparison과 결합된 경우
Load Comparison Inputs
→ Compare
→ Predicate
→ Load/Compute A and B
→ Select
→ Store Y
Mask를 memory에서 읽는 경우
Load Mask
→ Convert to Predicate
→ Load A
→ Load B
→ Select
→ Store Y
Pointer Selection
Load/Compute Addresses
→ Select Address
→ Load Selected Value
→ Store Y
36. 비용 구조
Tensor-Tensor Selection에서 predicate와 두 후보를 모두 memory에서 읽는다고 하자.
FP32 후보, byte mask 기준:
mask load: 약 1 byte 이상
a[i] load: 4 bytes
b[i] load: 4 bytes
y[i] store: 4 bytes
원소당 memory traffic은 최소 약 13 bytes 이상이 될 수 있다.
Mask가 4-byte integer로 저장되면 더 커질 수 있다.
반면 predicate를 같은 kernel에서 생성하면 mask load를 제거할 수 있다.
Compare input load
→ Predicate register
→ Selection
후보 중 하나가 scalar라면 후보 load도 줄어든다.
y[i] = p[i] ? x[i] : 0
주요 traffic:
x[i] load
y[i] store
Predicate가 local하게 생성된다면 매우 단순한 구조가 된다.
37. 표현 계층별 형태
37.1 수학 계층
y[i] =
a[i], if p[i]
b[i], otherwise
37.2 연산 그래프 계층
P ─┐
A ─┼→ Select → Y
B ─┘
일부 그래프 IR에서는 다음 이름을 사용할 수 있다.
- Select
- Where
- Conditional
- Masked Fill
- Choose
37.3 Loop 계층
for (int i = 0; i < n; ++i) {
y[i] = predicate[i] ? a[i] : b[i];
}
Comparison과 결합하면:
for (int i = 0; i < n; ++i) {
y[i] = x[i] > 0.0f ? x[i] : 0.0f;
}
37.4 CUDA Kernel 계층
__global__ void select_f32(
const bool* predicate,
const float* a,
const float* b,
float* y,
int n)
{
int i =
blockIdx.x * blockDim.x
+ threadIdx.x;
if (i < n) {
y[i] =
predicate[i] ? a[i] : b[i];
}
}
Compare-Select fusion:
if (i < n) {
bool p = x[i] > 0.0f;
y[i] = p ? x[i] : 0.0f;
}
37.5 PTX 계층
개념적으로 다음 구조가 나타날 수 있다.
setp
selp
예:
setp.gt.f32 predicate, x, 0
selp.f32 y, x, 0, predicate
정확한 syntax와 type modifier는 PTX version에 따라 달라질 수 있다.
37.6 SASS 계층
Architecture에 따라 다음 종류의 명령을 예상할 수 있다.
FSETP
ISETP
FSEL
SEL
PRMT 또는 bitwise mask 조합
Predicated MOV
정확한 mnemonic은 architecture와 compiler에 따라 달라질 수 있다.
핵심은 다음 데이터 흐름이다.
Compare
→ Predicate
→ Candidate Selection
38. SASS에서 Selection을 식별하는 기준
38.1 Predicate 입력
Selection은 predicate 또는 condition code에 의존한다.
P0 ───────────┐
R_true ───────┼→ Select → R_out
R_false ──────┘
38.2 두 후보 source
일반적인 Selection은 true와 false 후보를 가진다.
38.3 하나의 destination
선택된 값 하나가 destination register에 기록된다.
38.4 명시적 Branch 부재 가능성
Branchless Selection에서는 control-flow branch 없이 predicate 기반 value selection이 나타난다.
38.5 Predicate가 바로 소비됨
Comparison 결과가 별도의 일반 register나 memory에 저장되지 않고 즉시 Selection에 사용될 수 있다.
38.6 Direct Min/Max로 대체 가능
고수준 Selection이 SASS에서는 직접 maximum/minimum instruction으로 나타날 수 있다.
39. Selection이 제거되는 경우
동일 후보
select(p, a, a)
→ a
Constant Predicate
select(true, a, b)
→ a
select(false, a, b)
→ b
Predicate가 입력 범위로 결정됨
Compiler range analysis로 predicate가 항상 true 또는 false임을 증명할 수 있다.
Redundant Nested Selection
select(p, select(p, a, b), c)
에서 predicate 관계를 이용해 단순화할 수 있다.
Identity Candidate
특정 함수와 조합하면 불필요한 Selection을 제거할 수 있다.
40. Selection이 Branch로 변환될 수 있는 경우
Selection 후보 계산이 비싸고 predicate에 따라 한쪽만 필요하다면 compiler가 branch 구조를 선택할 수 있다.
예:
y =
p ? expensive_f(x)
: expensive_g(x)
다음 요소가 판단에 영향을 준다.
- 후보 계산 비용
- Predicate 분포
- Warp divergence 가능성
- Register pressure
- Memory load 비용
- Side effect
- Instruction count
Selection syntax를 사용했다고 해서 반드시 branchless instruction이 생성되는 것은 아니다.
41. Selection과 Register Pressure
두 후보를 모두 계산한 뒤 Selection하면 두 후보가 동시에 살아 있어야 할 수 있다.
R_a = compute_a(...)
R_b = compute_b(...)
R_y = select(P, R_a, R_b)
이 경우:
R_a lifetime
R_b lifetime
Predicate lifetime
이 겹치면서 register pressure가 증가할 수 있다.
Branch에서는 한 경로의 값만 살아 있을 가능성이 있다.
따라서 branchless Selection은 divergence를 줄이는 대신 register 사용량을 늘릴 수 있다.
42. Selection과 Instruction-Level Parallelism
두 후보 계산이 독립적이라면 동시에 진행할 수 있다.
a = f(x)
b = g(x)
이후:
y = select(p, a, b)
후보 계산 사이에는 dependency가 없으므로 ILP를 활용할 수 있다.
그러나 최종 Selection은 두 후보가 모두 준비될 때까지 기다려야 한다.
f 완료 ─┐
├→ Select
g 완료 ─┘
한 후보의 latency가 길면 critical path가 늘어날 수 있다.
43. Selection과 Fusion
Selection은 앞뒤 연산과 fusion하기 쉽다.
예:
p[i] = x[i] > 0
t[i] = select(p[i], x[i], 0)
y[i] = scale[i]*t[i] + bias[i]
분리 구조:
Compare
→ mask store
mask load
→ Select
→ t store
t load
→ FMA
→ y store
융합 구조:
Load X
→ Compare
→ Select
→ FMA
→ Store Y
중간 mask와 t materialization을 제거할 수 있다.
하지만 fusion으로 다음이 증가할 수 있다.
- Register pressure
- Instruction count
- Dependency chain
- Code size
- 후보 계산 중복
44. 허용되는 변환
Predicate 반전과 후보 교환
select(p, a, b)
→
select(not p, b, a)
동일 후보 제거
select(p, a, a)
→ a
Constant Predicate Folding
select(true, a, b)
→ a
Comparison-Selection Fusion
Predicate mask를 materialize하지 않고 바로 소비한다.
Branchless 변환
짧고 순수한 조건부 값을 Selection으로 바꿀 수 있다.
Direct Min/Max 변환
특수값 semantics가 일치할 때 적용할 수 있다.
Common Operation Hoisting
두 후보에 공통으로 적용되는 연산을 Selection 밖으로 이동할 수 있다.
Pointer Selection
두 값을 모두 load하는 대신 address를 먼저 선택할 수 있다.
Kernel Fusion
Comparison, Selection, 후속 Elementwise 연산을 결합할 수 있다.
45. 제한되는 변환
비선택 후보의 Side Effect 무시
두 후보를 모두 계산하면 원래 branch 의미와 달라질 수 있다.
정의되지 않은 후보의 선계산
Division by zero, invalid memory access 등이 발생할 수 있다.
Arithmetic Mask와 완전 동일시
NaN, Inf와 signed zero에서 결과가 달라질 수 있다.
Direct Min/Max와 완전 동일시
NaN 전파 규칙이 다를 수 있다.
후보 dtype 변환 무시
Type promotion과 rounding이 결과에 영향을 줄 수 있다.
Predicate의 NaN 의미 무시
Predicate가 floating-point comparison에서 왔다면 ordered/unordered semantics를 보존해야 한다.
Branchless가 항상 빠르다고 가정
후보 계산이 비싸거나 memory load가 많으면 더 느릴 수 있다.
46. 정확성 검증
기본 reference는 다음과 같다.
reference[i] =
a[i], if predicate[i] is true
b[i], otherwise
검증 항목:
- True predicate
- False predicate
- Alternating predicate
- Warp 전체 true
- Warp 전체 false
- Warp 내부 혼합
- Scalar candidate
- Tensor candidate
- Broadcasting
- NaN candidate
- Inf candidate
- Signed zero
- Mixed dtype
- Aliasing
- In-place update
수치 출력이라면 다음을 측정한다.
bitwise equality
max_abs_diff
max_relative_diff
논리적 선택 오류는 다음으로 측정할 수 있다.
mismatch_count
first_mismatch_index
47. 중요한 입력 패턴
Predicate 분포
all true
all false
50/50 random
alternating
warp별 uniform
warp 내부 random
후보 값
일반 유한값
NaN
+Inf
-Inf
+0.0
-0.0
큰 값
작은 값
후보 계산 비용
단순 scalar
memory load
FMA
exp
division
복잡한 polynomial
주소 선택
두 후보 모두 유효
한 후보만 유효
비연속 주소
coalesced 주소
이 입력들은 Selection과 Branch의 실제 차이를 드러내는 데 중요하다.
48. 실험 설계
실험 1: Predicate Tensor Selection
y[i] = predicate[i] ? a[i] : b[i]
목적:
- Predicate load
- 후보 두 개 load
- Selection instruction
- Output store 구조 확인
실험 2: Compare-Select Fusion
y[i] = x[i] > 0 ? x[i] : 0
목적:
- Comparison predicate의 즉시 소비
- Mask materialization 제거
- ReLU lowering 기반 마련
실험 3: Explicit Branch
if x[i] > 0:
y[i] = x[i]
else:
y[i] = 0
목적:
- Branch와 Selection SASS 비교
- Compiler가 branchless 형태로 변환하는지 확인
실험 4: Expensive Candidate Selection
y[i] =
p[i] ? exp(a[i]) : reciprocal(b[i])
목적:
- 두 후보 모두 계산되는지 확인
- Branch 생성 여부
- Predicate 분포에 따른 성능 비교
실험 5: Scalar False Candidate
y[i] = p[i] ? x[i] : 0
목적:
- Scalar/immediate 처리
- Tensor-Tensor Selection과 memory traffic 비교
실험 6: Mask Materialization
mask[i] = x[i] > 0
y[i] = mask[i] ? a[i] : b[i]
두 kernel로 분리한다.
목적:
- Mask global store/load 비용
- Fused Compare-Select와 비교
실험 7: Arithmetic Mask
m[i] = p[i] ? 1.0f : 0.0f
y[i] = m[i]*a[i] + (1-m[i])*b[i]
목적:
- Selection instruction과 arithmetic mask 비교
- NaN/Inf 특수값 차이
- Instruction count 비교
실험 8: Direct Maximum
다음을 비교한다.
select(x > z, x, z)
max(x, z)
목적:
- SASS 차이
- NaN 처리
- Signed zero 처리
- Performance 비교
실험 9: Pointer Selection
ptr = p ? ptr_a : ptr_b
y = *ptr
목적:
- Address selection
- 두 후보 load 방식과 비교
- Memory coalescing 분석
실험 10: In-place Conditional Update
state[i] =
update[i] ? candidate[i] : state[i]
목적:
- 기존 output을 false 후보로 사용하는 구조
- Alias와 register lifetime 확인
실험 11: Warp Predicate Distribution
다음 분포를 비교한다.
all true
all false
half true
alternating
random
목적:
- Selection과 Branch 성능 차이
- Warp divergence 영향
- Predication 비용 분석
실험 12: Attention Masked Fill
score[i, j] =
j <= i ? score[i, j] : -Inf
목적:
- Index comparison
- Selection
- Mask tensor 제거
- Attention score와 fusion 가능성 확인
49. 분석 항목
의미 수준
- Predicate 의미
- True/false 후보
- Broadcasting domain
- 후보 dtype
- Side effect 여부
- 후보 계산의 정의 가능성
코드 수준
- Selection syntax
- Explicit branch
- Candidate precomputation
- Mask materialization
- Pointer selection
- In-place 여부
PTX 수준
- Set predicate
- Select instruction
- Predicate branch
- Type conversion
- Predicated load/store
SASS 수준
- Predicate 생성
- FSEL 또는 SEL
- Predicated move
- Branch instruction
- 후보 load 수
- Register lifetime
- Mask load/store
실행 수준
- Kernel time
- Branch efficiency
- Warp divergence
- Register count
- Instruction count
- Memory throughput
- Predicate 분포별 성능
- Candidate 비용별 성능
정확성 수준
- Candidate 선택 일치
- NaN/Inf
- Signed zero
- Invalid candidate 평가 여부
- Arithmetic mask 차이
- Min/max 차이
50. 연구에서의 의미
Selection은 Comparison에서 생성된 논리적 관계를 실제 데이터 흐름으로 변환하는 primitive다.
핵심 관계는 다음과 같다.
Comparison
→ Predicate
→ Selection
→ Output Value
Selection을 통해 다음 개념을 연구할 수 있다.
논리적 중간값의 local 유지
Predicate를 global mask로 저장하지 않고
register에서 바로 소비할 수 있다.
Control flow의 data flow 변환
짧은 Branch
→ Predicate 기반 Selection
비선택 후보의 비관측성
수학적으로 비선택 후보는 출력에 영향을 주지 않지만, 실제 실행에서는 계산되거나 load될 수 있다.
Branchless 실행의 한계
Branch를 제거해도 후보 계산 비용과 register pressure가 증가할 수 있다.
특수값과 의미 보존
Arithmetic mask, min/max, direct Selection은 일반 유한값에서는 같아 보여도 NaN과 Inf에서 달라질 수 있다.
Fusion 가능성
Comparison, Selection, 후속 Elementwise 연산을 하나의 실행 모티프로 결합할 수 있다.
51. 최종 정의
Selection은 다음과 같이 정의할 수 있다.
Predicate의 참·거짓 값에 따라
두 후보 중 하나를 동일한 출력 위치의 값으로 선택하는 연산
기본 수식:
y[i] =
a[i], if p[i] is true
b[i], otherwise
또는:
y[i] = select(p[i], a[i], b[i])
핵심 의미:
1. Predicate에 의한 후보 선택
2. True 후보와 false 후보의 구분
3. 출력 원소 사이의 독립성
4. 비선택 후보는 수학적 출력에 영향을 주지 않음
5. Branch 없이 data flow로 표현 가능
기본 실행 구조:
Predicate
→ Candidate A
→ Candidate B
→ Select
→ Output
Comparison과 결합하면:
Load Inputs
→ Compare
→ Predicate
→ Select
→ Store
대표적인 최적화 가능성:
mask materialization 제거
comparison-selection fusion
branchless execution
predicate reuse
constant folding
동일 후보 제거
common operation hoisting
pointer selection
kernel fusion
중요한 제한:
Selection이 branchless라고 해서
선택되지 않은 후보의 계산 비용이 자동으로 사라지는 것은 아니다.
후보 계산이 side effect를 가지거나
정의되지 않은 연산 또는 invalid memory access를 포함하면
Branch를 Selection으로 단순 변환할 수 없다.
'SASS_Probe' 카테고리의 다른 글
| 기본 연산 의미 명세 08 - ReLU (1) | 2026.06.23 |
|---|---|
| 기본 연산 의미 명세 07 - Maximum과 Minimum (0) | 2026.06.23 |
| 기본 연산 의미 명세 05 - Comparison (0) | 2026.06.22 |
| 기본 연산 의미 명세 04 - Fused Multiply Add, FMA (0) | 2026.06.22 |
| 기본 연산 의미 명세 03 - Elementwise Multiply (0) | 2026.06.22 |