relu_f32 분석

relu_f32 분석

문서 위치 추천:

notes/02_control/relu_f32.md

1. CUDA 코드

대상 커널:

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

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

high-level 의미:

y[i] = max(x[i], 0.0f);

---

2. SASS 핵심 구조

/*0010*/ S2R R4, SR_CTAID.X ;
/*0020*/ S2R R3, SR_TID.X ;
/*0030*/ IMAD R4, R4, c[0x0][0x0], R3 ;
/*0040*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT ;
/*0050*/ @P0 EXIT ;

/*0060*/ MOV R5, 0x4 ;
/*0080*/ IMAD.WIDE R2, R4, R5, c[0x0][0x160] ;
/*0090*/ LDG.E.CONSTANT R2, [R2.64] ;

/*00a0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x168] ;
/*00b0*/ FMNMX.FTZ R7, RZ, R2, !PT ;
/*00c0*/ STG.E [R4.64], R7 ;
/*00d0*/ EXIT ;

역할로 풀면:

R4 = blockIdx.x;
R3 = threadIdx.x;

R4 = R4 * blockDim.x + R3;   // i

if (R4 >= n) {
    exit;
}

R5 = 4;                      // sizeof(float)

R2 = x_base + R4 * 4;         // &x[i]
R2 = load x[i];               // x[i]

R4 = y_base + R4 * 4;         // &y[i]

R7 = max(R2, 0.0f);           // ReLU

store y[i] = R7;

---

3. 명령어별 매핑

S2R R4, SR_CTAID.X	block index 읽기	blockIdx.x
S2R R3, SR_TID.X	thread index 읽기	threadIdx.x
IMAD R4, R4, c[0x0][0x0], R3	linear index 계산	i = blockIdx.x * blockDim.x + threadIdx.x
ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT	i >= n 비교	bounds check
@P0 EXIT	범위 밖 thread 종료	if (i >= n) return
MOV R5, 0x4	float byte stride	sizeof(float)
IMAD.WIDE R2, R4, R5, c[0x0][0x160]	x[i] 주소 계산	&x[i]
LDG.E.CONSTANT R2, [R2.64]	x[i] load	float v = x[i]
IMAD.WIDE R4, R4, R5, c[0x0][0x168]	y[i] 주소 계산	&y[i]
FMNMX.FTZ R7, RZ, R2, !PT	max(0, x[i])	v > 0 ? v : 0
STG.E [R4.64], R7	결과 저장	y[i] = relu(v)

 

---

4. 핵심 명령어: FMNMX.FTZ

FMNMX.FTZ R7, RZ, R2, !PT ;

FMNMX는 floating-point min/max 계열 명령어입니다.

여기서:

RZ = zero register
R2 = x[i]

따라서 피연산자는:

0.0f
x[i]

입니다.

그리고 결과는 R7에 들어갑니다.

R7 = max(0.0f, x[i])

즉 ReLU입니다.

---

5. RZ 의미

RZ는 zero register입니다.

RZ = 0

일반 레지스터처럼 보이지만 항상 0으로 읽힙니다.

그래서:

FMNMX.FTZ R7, RZ, R2, !PT ;

는 high-level로:

R7 = max(0.0f, R2);

로 보면 됩니다.

---

6. 왜 branch가 없는가?

CUDA source는 조건식입니다.

y[i] = v > 0.0f ? v : 0.0f;

그래서 naive하게 생각하면 이런 SASS를 예상할 수 있습니다.

FSETP
@P BRA
MOV

또는:

FSETP
SEL

하지만 실제 SASS는:

FMNMX.FTZ

입니다.

즉 컴파일러가 이 조건식을 다음 형태로 인식한 겁니다.

max(v, 0.0f)

그리고 하드웨어 primitive인 min/max 명령어로 낮췄습니다.

이건 매우 중요한 관찰입니다.

source-level branch expression이
SASS-level branch로 내려가지 않을 수 있다.

ReLU는 조건문처럼 쓰였지만 실제로는 branchless max operation이 되었습니다.

---

7. add/mul/fma와 비교

지금까지 산술/제어 패턴은 이렇게 정리됩니다.

CUDA expression              SASS primitive
------------------------------------------------
a + b                        FADD.FTZ
a * b                        FMUL.FTZ
a * b + c                    FFMA.FTZ
v > 0 ? v : 0                FMNMX.FTZ

즉 ReLU도 복잡한 control flow가 아니라 단일 arithmetic/select primitive로 내려갔습니다.

이건 NN operator 분석에서 중요합니다.

ReLU layer는 SASS에서 별도의 branch 구조가 아니라,
대부분 LDG → FMNMX → STG 패턴으로 관찰될 수 있다.

---

8. copy_global과 비교

copy_global:

LDG.E.CONSTANT R3, [R2.64] ;
STG.E [R4.64], R3 ;

relu_f32:

LDG.E.CONSTANT R2, [R2.64] ;
FMNMX.FTZ R7, RZ, R2, !PT ;
STG.E [R4.64], R7 ;

차이는 명확합니다.

copy_global:
    load x[i]
    store y[i]

relu_f32:
    load x[i]
    max(x[i], 0)
    store y[i]

즉 ReLU는 copy_global의 memory skeleton에 FMNMX 하나가 삽입된 형태입니다.

---

9. 레지스터 흐름

R4
 ├─ blockIdx.x
 ├─ i
 └─ y[i] address

R3
 ├─ threadIdx.x

R2
 ├─ x[i] address
 └─ x[i] value

R5
 └─ 4, sizeof(float)

R7
 └─ relu(x[i]) = max(x[i], 0)

특히 R2는 주소였다가 값으로 덮입니다.

IMAD.WIDE R2, R4, R5, c[0x0][0x160] ;
LDG.E.CONSTANT R2, [R2.64] ;

즉:

R2 = &x[i]
R2 = x[i]

이 패턴은 앞에서 본 copy_global, add_f32, mul_f32와 같습니다.

---

10. 패턴 문서에 추가

notes/patterns/control_f32.md를 새로 만들고 아래 내용을 추가하세요.

# Pattern: F32 ReLU

## CUDA

```cpp
float y = x > 0.0f ? x : 0.0f;

또는:

float y = max(x, 0.0f);

Observed SASS

FMNMX.FTZ R_out, RZ, R_x, !PT ;

Meaning

R_out = max(0.0f, R_x);

Observed in

• kernels/02_control/relu_f32.cu

Notes

• source-level ternary expression이 branch로 내려가지 않았다.
• FSETP + SEL이 아니라 FMNMX 단일 명령어로 표현되었다.
• RZ는 zero register로, ReLU의 threshold 0.0f 역할을 한다.
• ReLU는 SASS에서 LDG → FMNMX → STG 패턴으로 나타날 수 있다.

---

## 11. 결론

`relu_f32` 실험의 결론은 이겁니다.

```text
ReLU는 threshold operation이지만,
SASS에서는 branch instruction이 아니라 FMNMX instruction으로 표현되었다.

더 정확히는:

v > 0 ? v : 0
→ compiler가 max(v, 0)로 rewrite
→ SASS에서 FMNMX.FTZ로 lowering

따라서 현재 분석 프레임에 넣으면 이렇게 됩니다.

CUDA source expression:
    v > 0 ? v : 0

compiler graph rewrite:
    max(v, 0)

SASS primitive:
    FMNMX.FTZ

intermediate materialization:
    없음. load한 값이 바로 FMNMX를 거쳐 store됨.

이게 핵심입니다.

다음은 clamp_f32.sass를 보면 됩니다.
거기서는 min(max(x, lo), hi)가 FMNMX 두 개로 내려가는지, 아니면 FSETP + SEL로 내려가는지가 포인트입니다.