기본 연산 의미 명세 08 - ReLU

1. 연산 개요

ReLU는 입력이 양수이면 그대로 유지하고, 음수이면 0으로 바꾸는 활성화 함수다.

ReLU는 Rectified Linear Unit의 약자다.

기본 정의는 다음과 같다.

y[i] = max(x[i], 0)

구간별로 표현하면 다음과 같다.

x[i] > 0이면:
    y[i] = x[i]

x[i] <= 0이면:
    y[i] = 0

ReLU는 다음 세 기본 연산 관점에서 해석할 수 있다.

Maximum 관점:
y[i] = max(x[i], 0)

Comparison과 Selection 관점:
p[i] = x[i] > 0
y[i] = select(p[i], x[i], 0)

Control-flow 관점:
if x[i] > 0:
    y[i] = x[i]
else:
    y[i] = 0

수학적 정의는 같아 보이지만 실제 부동소수점 실행에서는 NaN과 signed zero 처리에 따라 결과가 달라질 수 있다.

---

2. 수학적 정의

ReLU 함수는 다음과 같이 정의한다.

ReLU(x) = max(0, x)

Piecewise 형태는 다음과 같다.

ReLU(x) =
    x, if x > 0
    0, otherwise

Elementwise tensor 연산에서는 다음과 같다.

y[i] = ReLU(x[i])

입력과 출력 shape은 동일하다.

shape(Y) = shape(X)

입력 원소 하나가 출력 원소 하나에 대응한다.

x[i]
→ ReLU
→ y[i]

---

3. ReLU의 기하학적 의미

ReLU는 실수축을 두 영역으로 나눈다.

음수 영역:
x < 0
→ 출력 0

양수 영역:
x > 0
→ 입력 그대로 통과

그래프는 다음 특징을 가진다.

x < 0:
기울기 0

x > 0:
기울기 1

입력과 출력의 관계는 다음과 같다.

x축 음수 영역은 0으로 눌림
x축 양수 영역은 항등적으로 유지

따라서 ReLU는 음수 영역을 제거하는 반면 양수 영역에서는 선형성을 유지한다.

---

4. 입력과 출력 domain

일반적인 ReLU는 부동소수점 tensor에 적용된다.

주요 dtype은 다음과 같다.

FP32
FP16
BF16
FP64
FP8 계열

정수 tensor에도 적용할 수 있다.

y[i] = max(x[i], 0)

정수에서는 NaN과 signed zero가 없으므로 의미가 더 단순하다.

출력은 항상 입력과 같은 범위에 있지 않을 수 있지만 다음 관계를 만족한다.

y[i] >= 0

유한한 입력에 대해서는:

0 <= y[i] <= max(x[i], 0)

실제로는 정의 자체가:

y[i] = max(x[i], 0)

이므로 양수 출력은 원래 입력값과 같다.

---

5. 데이터 의존성

ReLU는 단항 Elementwise Map이다.

y[i] = f(x[i])

여기서:

f(x) = max(x, 0)

출력 y[i]는 대응하는 입력 x[i]에만 의존한다.

x[0] → y[0]
x[1] → y[1]
x[2] → y[2]

다른 위치의 입력에는 의존하지 않는다.

i != j이면
y[i]는 x[j]를 필요로 하지 않는다.

따라서 ReLU 자체에는 다음 구조가 필요하지 않다.

Reduction
Warp shuffle
Shared memory communication
Block synchronization
Atomic operation

---

6. 원소별 독립성

모든 출력 원소를 독립적으로 계산할 수 있다.

thread 0 → y[0]
thread 1 → y[1]
thread 2 → y[2]
...

출력 계산 순서는 의미에 영향을 주지 않는다.

y[0]을 먼저 계산해도 됨
y[2]를 먼저 계산해도 됨
모두 동시에 계산해도 됨

이 성질은 ReLU를 GPU에서 쉽게 병렬화할 수 있게 한다.

---

7. 핵심 의미 불변성

ReLU 구현이 달라져도 다음 조건은 유지되어야 한다.

7.1 음수 제거

유한한 음수 입력에 대해서:

x < 0이면
ReLU(x) = 0

7.2 양수 보존

유한한 양수 입력에 대해서:

x > 0이면
ReLU(x) = x

7.3 비음수 출력

일반적인 유한 입력에 대해서:

ReLU(x) >= 0

7.4 위치 대응

각 출력은 같은 위치의 입력으로부터 계산되어야 한다.

y[i] = ReLU(x[i])

7.5 Shape 보존

shape(Y) = shape(X)

---

8. ReLU의 대수적 성질

8.1 멱등성

ReLU를 두 번 적용해도 결과는 한 번 적용한 것과 같다.

ReLU(ReLU(x)) = ReLU(x)

첫 번째 ReLU 출력은 이미 음수가 아니기 때문이다.

이 성질은 중복 ReLU 제거의 근거가 된다.

ReLU
→ ReLU

를:

ReLU

하나로 줄일 수 있다.

---

8.2 양의 동차성

0 이상의 scalar a에 대해 다음이 성립한다.

ReLU(a*x) = a*ReLU(x)

단:

a >= 0

이어야 한다.

예를 들어:

ReLU(2*x) = 2*ReLU(x)

하지만 a < 0이면 일반적으로 성립하지 않는다.

ReLU(-x)
!=
-ReLU(x)

---

8.3 단조성

입력이 커지면 출력도 감소하지 않는다.

x <= z이면
ReLU(x) <= ReLU(z)

따라서 ReLU는 단조 증가 함수다.

---

8.4 비선형성

ReLU는 전체 실수 범위에서 선형 함수가 아니다.

일반적으로:

ReLU(x + z)
!=
ReLU(x) + ReLU(z)

예:

x = 1
z = -1

이면:

ReLU(x + z) = ReLU(0) = 0

하지만:

ReLU(x) + ReLU(z) = 1 + 0 = 1

이다.

---

8.5 Convexity

ReLU는 convex 함수다.

즉 두 값 x, z와 0 <= t <= 1에 대해:

ReLU(t*x + (1-t)*z)
<=
t*ReLU(x) + (1-t)*ReLU(z)

가 성립한다.

이 성질은 최적화 이론에서는 중요하지만, 직접적인 SASS lowering 구조를 결정하는 핵심 요소는 아니다.

---

9. ReLU와 Maximum

ReLU는 scalar Maximum이다.

ReLU(x) = max(x, 0)

따라서 direct maximum instruction을 사용할 수 있다.

추상 실행 구조:

Load X
→ Maximum with 0
→ Store Y

SASS 수준에서는 architecture에 따라 다음과 같은 max 계열 명령으로 나타날 수 있다.

FMAX
FMNMX

또는 유사한 floating-point maximum 명령이 사용될 수 있다.

---

10. ReLU와 Comparison-Selection

ReLU를 다음처럼 분해할 수 있다.

p = x > 0
y = select(p, x, 0)

실행 구조:

Load X
→ Compare X with 0
→ Predicate 생성
→ Select X or 0
→ Store Y

PTX 또는 SASS에서는 다음과 같은 패턴이 나타날 수 있다.

Comparison
→ Predicate
→ Selection

개념적으로:

FSETP
→ FSEL

또는 이에 대응하는 명령 조합이 사용될 수 있다.

---

11. ReLU와 Branch

조건문을 직접 사용하면 다음과 같다.

if x > 0:
    y = x
else:
    y = 0

실행 구조:

Compare
→ Predicate
→ Branch
→ Move X 또는 Zero

GPU에서는 warp 안의 thread별로 부호가 다르면 divergence가 발생할 수 있다.

예:

thread 0: x > 0
thread 1: x <= 0
thread 2: x > 0
thread 3: x <= 0

이 경우 true와 false 경로를 나누어 실행할 수 있다.

그러나 ReLU의 두 경로는 매우 짧기 때문에 compiler는 branch보다 max 또는 selection을 선택할 가능성이 높다.

실제 lowering은 SASS로 확인해야 한다.

---

12. Branchless ReLU

Branchless ReLU는 explicit control-flow branch 없이 실행된다.

대표적인 형태는 다음과 같다.

y = max(x, 0)

또는:

p = x > 0
y = select(p, x, 0)

장점:

Warp divergence 제거 가능
짧은 dependency chain
예측 가능한 instruction flow

ReLU처럼 후보 계산이 단순한 경우 branchless 표현이 자연스럽다.

---

13. Arithmetic Mask ReLU

ReLU를 boolean mask와 Multiply로 표현할 수도 있다.

mask = x > 0 ? 1 : 0
y = x * mask

유한한 일반 입력에서는 ReLU와 같은 결과를 만들 수 있다.

하지만 다음 문제가 있다.

추가 연산

Comparison
Predicate-to-number 변환
Multiply

특수값 문제

x = +Inf이고 mask가 1이면 문제없다.

하지만 x = NaN이면:

NaN > 0 → false
mask = 0
NaN * 0 = NaN

Selection 방식에서는 구현에 따라 0을 선택할 수도 있다.

따라서 arithmetic mask와 direct maximum은 특수값에서 같은 의미를 보장하지 않는다.

---

14. NaN 처리

ReLU의 NaN 처리 방식은 구현 정의와 사용한 primitive에 따라 달라질 수 있다.

다음 입력을 생각하자.

x = NaN

Comparison-Selection

p = NaN > 0

일반적으로 false다.

따라서:

select(false, NaN, 0)

이면 0을 반환할 수 있다.

NaN 전파형 Maximum

max_propagate(NaN, 0) = NaN

숫자 우선형 Maximum

max_number(NaN, 0) = 0

따라서 ReLU 구현을 비교할 때 반드시 다음을 확인해야 한다.

NaN을 전파하는가
NaN을 0으로 바꾸는가
Operand 순서에 따라 결과가 달라지는가

---

15. Signed Zero

Floating-point에는 +0.0과 -0.0이 있다.

ReLU에서 입력이 -0.0일 때 어떤 결과가 나오는지 구현에 따라 달라질 수 있다.

Piecewise 정의

x > 0이면 x
그렇지 않으면 0

-0.0 > 0은 false이므로 false 후보의 +0.0을 선택할 수 있다.

ReLU(-0.0) = +0.0

Direct Maximum

Maximum instruction의 signed-zero 규칙에 따라:

max(-0.0, +0.0) = +0.0

가 기대될 수 있다.

하지만 operand 순서와 instruction semantics를 확인해야 한다.

Signed zero가 모델 정확도에는 거의 영향을 주지 않을 수 있지만 bitwise equivalence에서는 중요하다.

---

16. Infinity 처리

Positive Infinity

ReLU(+Inf) = +Inf

Negative Infinity

ReLU(-Inf) = 0

일반적인 maximum 정의에서는:

max(+Inf, 0) = +Inf
max(-Inf, 0) = 0

이다.

Infinity 처리 역시 NaN과 달리 정상적인 ordering에 참여한다.

---

17. Subnormal과 작은 음수

매우 작은 음수는 ReLU에서 0으로 제거된다.

x < 0이면
y = 0

매우 작은 양수는 그대로 유지되어야 한다.

x > 0이면
y = x

하지만 GPU가 subnormal을 flush-to-zero하는 설정이라면 매우 작은 양수가 0으로 바뀔 수 있다.

따라서 다음을 구분해야 한다.

ReLU가 0으로 만든 것

과:

입력 또는 실행 환경이 subnormal을 0으로 처리한 것

---

18. ReLU의 출력 범위

유한 입력에 대해 다음이 성립한다.

y >= 0

또한:

x >= 0이면 y = x
x < 0이면 y = 0

출력 범위는 다음과 같다.

Y ∈ [0, +Inf)

NaN을 전파하는 구현에서는 NaN이 출력에 포함될 수 있다.

---

19. ReLU의 미분

ReLU의 미분은 다음과 같다.

x > 0이면:
    dReLU/dx = 1

x < 0이면:
    dReLU/dx = 0

x = 0에서는 고전적인 의미의 미분이 정의되지 않는다.

딥러닝 구현에서는 일반적으로 특정 subgradient를 선택한다.

대개:

x = 0이면 gradient = 0

로 처리한다.

하지만 정확한 convention은 구현별로 명세해야 한다.

---

20. ReLU Backward

Forward:

y[i] = max(x[i], 0)

Backward:

dx[i] =
    dy[i], if x[i] > 0
    0,     otherwise

Selection 형태:

p[i] = x[i] > 0
dx[i] = select(p[i], dy[i], 0)

또는 mask Multiply:

dx[i] = dy[i] * mask[i]

여기서:

mask[i] = 1, if x[i] > 0
mask[i] = 0, otherwise

---

21. Backward에서 필요한 정보

ReLU backward는 입력 x의 부호가 필요하다.

x[i] > 0인지 여부

하지만 forward 출력 y를 이용할 수도 있다.

유한 입력과 일반적인 ReLU에서는:

y[i] > 0

이면 원래 입력도 양수다.

따라서 다음과 같이 쓸 수 있다.

dx[i] =
    dy[i], if y[i] > 0
    0, otherwise

이 경우 원래 입력을 저장하지 않고 출력만으로 backward mask를 복원할 수 있다.

다만 다음을 확인해야 한다.

• NaN 처리
• Signed zero
• In-place 실행
• 정확한 x = 0 gradient convention
• Output이 이후 덮어써지는지 여부

---

22. Mask 저장과 재계산

ReLU backward를 위해 forward에서 mask를 저장할 수 있다.

mask[i] = x[i] > 0

Backward:

dx[i] = select(mask[i], dy[i], 0)

장점:

Backward에서 comparison 제거

단점:

Mask memory allocation
Mask global store
Mask global load

반대로 mask를 재계산할 수 있다.

mask[i] = x[i] > 0

또는:

mask[i] = y[i] > 0

이 선택은 다음 trade-off를 가진다.

저장:
추가 memory traffic

재계산:
추가 comparison

ReLU처럼 comparison이 저렴한 경우 재계산이 유리할 수 있다.

---

23. In-place ReLU

ReLU는 입력 buffer를 직접 갱신할 수 있다.

x[i] = max(x[i], 0)

장점:

별도 output buffer 불필요
메모리 사용량 감소

기본 memory traffic은 다음과 같다.

x load
x store

안전 조건은 다음과 같다.

원래 입력이 다른 consumer에서 필요하지 않음
Backward에서 원래 음수 값이 필요하지 않음
Alias 문제가 없음

ReLU backward는 보통 부호 정보만 필요하므로 in-place가 비교적 자연스럽다.

---

24. ReLU와 Kernel Fusion

ReLU는 앞선 연산의 epilogue에 쉽게 fusion할 수 있다.

예:

t[i] = a*x[i] + b
y[i] = ReLU(t[i])

분리 실행:

Load X
→ FMA
→ Store T

Load T
→ ReLU
→ Store Y

융합 실행:

Load X
→ FMA
→ ReLU
→ Store Y

제거되는 것은 다음과 같다.

T global store
T global load
추가 kernel launch
추가 index 계산

---

25. GEMM + Bias + ReLU

대표적인 epilogue fusion은 다음과 같다.

C = A * B
T = C + bias
Y = ReLU(T)

분리 구조:

GEMM
→ C global store

Bias Add kernel
→ T global store

ReLU kernel
→ Y global store

융합 구조:

GEMM register accumulator
→ Bias Add
→ ReLU
→ Final Store

SASS에서는 다음과 같은 산술 흐름을 기대할 수 있다.

FFMA accumulation
→ FADD 또는 FFMA bias
→ FMAX/FMNMX 또는 Select
→ STG

이 경우 ReLU 자체보다 중간 materialization 제거가 더 큰 최적화 효과를 만들 수 있다.

---

26. Convolution + ReLU

Convolution 결과에도 동일한 epilogue fusion을 적용할 수 있다.

acc = convolution accumulation
acc = acc + bias
y = max(acc, 0)

실행 구조:

Convolution accumulator
→ Bias
→ ReLU
→ Store

별도의 ReLU kernel을 실행하지 않으므로 memory traffic과 launch overhead를 줄일 수 있다.

---

27. Residual Add + ReLU

다음 구조를 생각할 수 있다.

t[i] = main[i] + residual[i]
y[i] = ReLU(t[i])

융합하면:

Load Main
→ Load Residual
→ Add
→ ReLU
→ Store Y

SASS에서는 다음과 같은 흐름이 가능하다.

LDG
LDG
FADD
FMAX 또는 Select
STG

Residual Add 결과가 다른 consumer에서 필요하지 않다면 중간 tensor materialization을 제거할 수 있다.

---

28. ReLU와 Sparsity

ReLU는 음수 입력을 0으로 바꾸므로 출력에 0이 많이 생길 수 있다.

negative input
→ zero output

이를 activation sparsity라고 볼 수 있다.

하지만 출력에 0이 많다고 해서 일반 dense GPU kernel이 자동으로 계산을 건너뛰는 것은 아니다.

다음 추가 구조가 필요할 수 있다.

• Sparse representation
• Zero detection
• Compaction
• Sparse kernel
• Structured sparsity
• Mask propagation

따라서 수학적 sparsity와 실제 실행량 감소는 별개의 문제다.

---

29. ReLU와 Memory-bound 특성

독립된 FP32 ReLU kernel에서 원소 하나당 주요 memory traffic은 다음과 같다.

x[i] load: 4 bytes
y[i] store: 4 bytes

합계: 약 8 bytes

산술은 maximum 또는 compare-select 정도다.

산술 연산: 매우 적음
Memory 이동: 상대적으로 큼

따라서 독립 ReLU kernel은 일반적으로 memory-bound가 되기 쉽다.

이것이 ReLU fusion이 중요한 이유다.

ReLU 계산 비용 자체보다
입출력 memory와 kernel launch 비용이 더 클 수 있음

---

30. 추상 실행 모티프

Direct Maximum ReLU

Index
→ Load X
→ Maximum with Zero
→ Store Y

Compare-Select ReLU

Index
→ Load X
→ Compare X > 0
→ Predicate
→ Select X or Zero
→ Store Y

Branch ReLU

Index
→ Load X
→ Compare
→ Branch
→ Move X or Zero
→ Store Y

Fused Epilogue ReLU

Producer Accumulator
→ Bias/Scale
→ ReLU
→ Final Store

---

31. 표현 계층별 형태

31.1 수학 계층

y[i] = max(x[i], 0)

31.2 연산 그래프 계층

X
→ ReLU
→ Y

또는 primitive로 분해하면:

X ─────────────┐
               ├→ Maximum → Y
Constant 0 ────┘

31.3 Loop 계층

for (int i = 0; i < n; ++i) {
    y[i] = x[i] > 0.0f ? x[i] : 0.0f;
}

또는:

for (int i = 0; i < n; ++i) {
    y[i] = fmaxf(x[i], 0.0f);
}

31.4 CUDA Kernel 계층

__global__ void relu_f32(
    const float* x,
    float* y,
    int n)
{
    int i =
        blockIdx.x * blockDim.x
        + threadIdx.x;

    if (i < n) {
        y[i] = fmaxf(x[i], 0.0f);
    }
}

Comparison-Selection 형태:

if (i < n) {
    float v = x[i];
    y[i] = v > 0.0f ? v : 0.0f;
}

31.5 PTX 계층

가능한 형태는 다음과 같다.

max.f32

또는:

setp.gt.f32
selp.f32

Boundary check를 위한 별도의 integer predicate도 나타날 수 있다.

31.6 SASS 계층

예상 가능한 구조는 다음과 같다.

S2R
→ Index/Address 계산
→ Boundary Predicate
→ LDG
→ FMAX/FMNMX
→ STG
→ EXIT

또는:

LDG
→ FSETP
→ FSEL
→ STG

정확한 mnemonic은 GPU architecture와 compiler version에 따라 달라질 수 있다.

---

32. SASS에서 ReLU를 식별하는 기준

32.1 하나의 입력 load

LDG X

32.2 Zero와의 비교 또는 maximum

MAX(X, 0)

또는:

Compare X > 0
→ Select X or 0

32.3 하나의 출력 store

STG Y

32.4 Cross-thread communication 부재

일반적으로 다음이 없다.

SHFL
BAR
Shared-memory reduction
Atomic operation

32.5 짧은 dependency chain

LDG
→ ReLU instruction
→ STG

32.6 Boundary predicate와 ReLU predicate 구분

Kernel boundary check에서도 predicate가 생성될 수 있다.

i < n

ReLU predicate는 데이터값에 대한 조건이다.

x > 0

SASS 분석에서는 두 predicate의 operand와 사용처를 구분해야 한다.

---

33. Direct Maximum이 생성되지 않는 경우

다음과 같은 이유로 Compare-Select가 생성될 수 있다.

NaN semantics 보존

Source 표현이 요구하는 NaN 처리와 direct maximum 명령 의미가 다를 수 있다.

Signed zero semantics

Tie-breaking과 zero 부호를 보존하기 위해 별도 comparison-selection이 필요할 수 있다.

Compiler 표현

Ternary expression이 그대로 predicate-select로 lowering될 수 있다.

데이터 타입

특정 dtype에서 direct max instruction 지원이 제한될 수 있다.

복합 조건

단순 ReLU가 아니라 추가 조건이 결합될 수 있다.

---

34. ReLU가 제거되는 경우

입력 범위가 항상 비음수가 아닌 경우가 아니라, 항상 비음수인 경우

Compiler나 graph optimizer가 다음을 증명할 수 있다고 하자.

x[i] >= 0

그러면:

ReLU(x[i]) = x[i]

이므로 ReLU를 제거할 수 있다.

중복 ReLU

ReLU(ReLU(x))
→ ReLU(x)

앞선 Clamp와 중복

Clamp(x, 0, upper)
→ ReLU

에서 ReLU가 불필요할 수 있다.

출력이 사용되지 않음

Dead code elimination으로 전체 ReLU를 제거할 수 있다.

---

35. ReLU의 연산 재배치

ReLU는 일반적으로 Add나 Multiply와 자유롭게 교환할 수 없다.

Add와의 관계

일반적으로:

ReLU(x + b)
!=
ReLU(x) + b

양의 Scale과의 관계

a >= 0이면:

ReLU(a*x) = a*ReLU(x)

따라서 양의 scalar scale은 ReLU 앞뒤로 이동할 수 있는 수학적 자유도가 있다.

하지만 부동소수점에서는 다음 차이를 확인해야 한다.

round(ReLU(round(a*x)))

과:

round(a*ReLU(x))

Overflow, underflow와 rounding이 달라질 수 있다.

음의 Scale

a < 0이면 이동할 수 없다.

ReLU(a*x)
!=
a*ReLU(x)

---

36. ReLU와 Bias 순서

다음 두 연산은 일반적으로 다르다.

ReLU(x + b)

ReLU(x) + b

예:

x = -1
b = 2

첫 번째:

ReLU(-1 + 2)
= ReLU(1)
= 1

두 번째:

ReLU(-1) + 2
= 0 + 2
= 2

따라서 GEMM + Bias + ReLU epilogue에서 연산 순서는 의미의 일부다.

GEMM
→ Bias Add
→ ReLU

와:

GEMM
→ ReLU
→ Bias Add

는 다른 연산이다.

---

37. ReLU6과의 구분

ReLU6는 다음과 같다.

ReLU6(x) = min(max(x, 0), 6)

즉 ReLU와 upper bound가 결합된 Clamp다.

ReLU6
=
Maximum with 0
+
Minimum with 6

일반 ReLU는 upper bound가 없다.

ReLU(x) = max(x, 0)

따라서 ReLU6는 별도의 연산으로 문서화해야 한다.

---

38. Leaky ReLU와의 구분

Leaky ReLU는 음수 영역을 완전히 0으로 만들지 않는다.

LeakyReLU(x) =
    x,       if x >= 0
    alpha*x, otherwise

Selection 형태:

y = select(x >= 0, x, alpha*x)

ReLU는 alpha = 0인 특수한 경우로 볼 수 있다.

ReLU(x)
=
LeakyReLU(x, alpha=0)

하지만 alpha*x 계산과 특수값 semantics 때문에 실제 lowering 구조는 다를 수 있다.

---

39. 허용되는 변환

Direct Maximum 사용

select(x > 0, x, 0)
→ max(x, 0)

NaN과 signed zero semantics가 일치해야 한다.

Branchless 변환

짧은 조건 branch를 maximum 또는 selection으로 바꿀 수 있다.

중복 ReLU 제거

ReLU(ReLU(x))
→ ReLU(x)

In-place 실행

x = ReLU(x)

Producer Fusion

Affine
→ ReLU

를 하나의 kernel로 결합할 수 있다.

GEMM/Convolution Epilogue Fusion

Accumulator에서 바로 ReLU를 적용할 수 있다.

Mask 재계산

Backward용 mask를 저장하지 않고 입력 또는 출력에서 다시 계산할 수 있다.

양의 Scale 이동

수학적 조건과 수치 오차가 허용되면 양의 scalar scale과 순서를 조정할 수 있다.

---

40. 제한되는 변환

NaN semantics 무시

Direct maximum과 comparison-selection이 NaN에서 다를 수 있다.

Signed zero 무시

-0.0 입력의 출력 부호가 구현에 따라 달라질 수 있다.

Bias와 순서 교환

ReLU(x+b)

와:

ReLU(x)+b

는 일반적으로 다르다.

음의 Scale 이동

a < 0

이면 ReLU와 scale을 교환할 수 없다.

Arithmetic Mask와 완전 동일시

x * (x > 0)

는 NaN과 Inf에서 다른 의미를 가질 수 있다.

Backward mask를 무조건 출력에서 복원

NaN 처리와 x = 0 convention을 확인해야 한다.

In-place 안전성 무시

원래 입력이 다른 consumer나 backward에서 필요할 수 있다.

---

41. 정확성 검증

기본 reference는 ReLU의 특수값 의미와 함께 정의해야 한다.

일반 유한값 reference:

reference[i] =
    x[i], if x[i] > 0
    0, otherwise

또는 direct maximum semantics:

reference[i] = specified_max(x[i], 0)

검증 항목은 다음과 같다.

일반 양수
일반 음수
0
+0.0
-0.0
NaN
+Inf
-Inf
Subnormal 양수
Subnormal 음수
최대 유한값
최소 normal 값

측정 항목:

bitwise equality
max_abs_diff
mismatch_count
NaN 여부
signed-zero bit

일반 ReLU는 유한 입력에서는 정확한 선택 연산이므로 정상값에서 수치 오차는 거의 없어야 한다.

---

42. ReLU Backward 정확성 검증

Backward reference:

dx[i] =
    dy[i], if x[i] > 0
    0, otherwise

검증해야 할 조건은 다음과 같다.

x > 0
x < 0
x = +0.0
x = -0.0
x = NaN
dy = NaN
dy = Inf

또한 다음 구현들을 비교할 수 있다.

입력 x로 mask 재계산
출력 y로 mask 재계산
Forward mask 저장
Arithmetic mask Multiply
Predicate Selection

---

43. 실험 설계

실험 1: Direct Maximum ReLU

y[i] = fmaxf(x[i], 0.0f)

목적:

Direct maximum instruction 확인
기본 LDG-MAX-STG 구조 확인
NaN 및 signed zero 결과 기록

실험 2: Ternary ReLU

y[i] =
    x[i] > 0.0f ? x[i] : 0.0f

목적:

Compare-Select lowering 확인
Direct maximum 구현과 SASS 비교

실험 3: Explicit Branch ReLU

if x[i] > 0:
    y[i] = x[i]
else:
    y[i] = 0

목적:

Compiler가 branch를 유지하는지 확인
Predication 또는 max로 변환하는지 확인

실험 4: Arithmetic Mask ReLU

mask = x[i] > 0 ? 1.0f : 0.0f
y[i] = x[i] * mask

목적:

Instruction 수 비교
NaN/Inf 결과 비교
Predicate-to-value 변환 확인

실험 5: In-place ReLU

x[i] = max(x[i], 0)

목적:

Out-of-place와 load/store 구조 비교
Memory 사용량 비교
Alias 조건 확인

실험 6: ReLU + Scale

y[i] = scale * max(x[i], 0)

목적:

ReLU 후 Multiply 구조 확인
양의 scale 재배치 가능성 검토
Fusion 효과 측정

실험 7: Affine + ReLU

y[i] = max(a*x[i] + b, 0)

목적:

FFMA + FMAX 구조 확인
분리 kernel과 fused kernel 비교

실험 8: Add + ReLU

y[i] = max(x[i] + z[i], 0)

목적:

FADD + FMAX
중간 materialization 제거
Residual Add + ReLU 구조 분석

실험 9: NaN/Signed Zero Matrix

입력:

NaN
+0.0
-0.0
+Inf
-Inf

구현 비교:

fmaxf
ternary
explicit branch
arithmetic mask

목적:

각 구현의 의미 차이 확인

실험 10: Backward Mask 저장

Forward에서 mask를 별도 tensor로 저장한다.

목적:

Mask memory 비용
Backward comparison 제거 효과

실험 11: Backward Mask 재계산

입력 또는 출력에서 mask를 재계산한다.

목적:

Memory와 computation trade-off 비교

실험 12: FP16/BF16 ReLU

목적:

Packed instruction 가능성
Input/output conversion
Subnormal 처리
FP32와 SASS 차이

---

44. 분석 항목

의미 수준

ReLU 정의
NaN 처리
Signed zero 처리
x = 0에서의 backward convention

코드 수준

Direct maximum
Ternary
Explicit branch
Arithmetic mask
In-place
Fusion

PTX 수준

max
setp
selp
mul
branch
type conversion

SASS 수준

FMAX/FMNMX 계열
FSETP
FSEL/SEL
BRA
LDG/STG
Predicate register

실행 수준

Kernel time
Effective memory bandwidth
Instruction count
Branch efficiency
Warp divergence
Register count
Launch overhead

정확성 수준

Bitwise equality
NaN output
Signed-zero output
Backward mismatch
Special-value matrix

---

45. 연구에서의 의미

ReLU는 지금까지 정의한 여러 primitive가 실제 하나의 활성화 연산으로 결합되는 첫 구체 사례다.

Elementwise Map
+
Comparison
+
Selection

또는:

Scalar Maximum

으로 해석할 수 있다.

ReLU를 통해 다음 관계를 관찰할 수 있다.

수학적 piecewise 정의
→ Comparison과 Selection

Comparison과 Selection
→ Direct Maximum instruction

Elementwise 독립성
→ Thread-level 병렬화

낮은 산술 집약도
→ Memory-bound kernel

Producer 결과의 즉시 소비
→ Epilogue Fusion

멱등성
→ 중복 ReLU 제거

출력 부호 정보
→ Backward mask 재구성 가능성

또한 ReLU는 “같은 정상값 결과를 만드는 구현들이 특수값에서도 같은가”라는 연구 질문을 명확하게 보여준다.

Direct maximum
Comparison-selection
Explicit branch
Arithmetic mask

는 일반 유한 입력에서는 같을 수 있지만 NaN과 signed zero에서는 달라질 수 있다.

---

46. 최종 정의

ReLU는 다음과 같이 정의할 수 있다.

입력이 양수이면 그대로 통과시키고,
양수가 아니면 0을 출력하는
원소별 비선형 활성화 연산

기본 수식:

y[i] = max(x[i], 0)

Piecewise 정의:

y[i] =
    x[i], if x[i] > 0
    0, otherwise

핵심 의미:

1. 음수 입력 제거
2. 양수 입력 보존
3. 비음수 출력
4. 출력 원소 간 독립성
5. Comparison-Selection 또는 Maximum으로 표현 가능

기본 실행 구조:

Index
→ Load X
→ Maximum with Zero
→ Store Y

또는:

Index
→ Load X
→ Compare X > 0
→ Select X or Zero
→ Store Y

대표적인 최적화 가능성:

Direct maximum instruction
Branchless lowering
In-place execution
중복 ReLU 제거
Affine + ReLU fusion
GEMM/Convolution epilogue fusion
Backward mask 재계산
Mask materialization 제거

중요한 제한:

Direct maximum,
comparison-selection,
explicit branch,
arithmetic mask는

일반적인 유한 입력에서는 같은 결과를 만들 수 있지만
NaN과 signed zero에서는 서로 다른 의미를 가질 수 있다.