SASS 명령어 분류, 내용 확인

좋아. 사실 SASS를 보기 시작하면 가장 먼저 해야 할 일은 명령어를 기능별로 분류하는 것이야.

처음에는 개별 명령어 의미를 외우기보다,

1. Memory
2. Arithmetic
3. Special Function
4. Data Movement
5. Control Flow
6. Tensor Core

로 나누는 게 훨씬 이해가 쉽다.

---

1. Memory 계열

GPU 메모리와 레지스터 사이의 데이터 이동

LDG

LDG.E R4, [R2]

의미

R4 = *R2;

즉

Global Memory
↓
Register

로드

---

예를 들어

float x = input[i];

를 컴파일하면

대부분

LDG

가 나온다.

---

STG

STG.E [R2], R4

의미

*R2 = R4;

즉

Register
↓
Global Memory

저장

---

예시

output[i] = x;

↓

STG

---

LDS

Shared Memory → Register

---

예시

float x = smem[tid];

↓

LDS

---

STS

Register → Shared Memory

---

예시

smem[tid] = x;

↓

STS

---

2. Arithmetic 계열

---

FADD

Float Add

FADD R1, R2, R3

의미

R1 = R2 + R3;

---

FMUL

Float Multiply

FMUL R1, R2, R3

의미

R1 = R2 * R3;

---

FFMA

Fused Multiply Add

FFMA R1, R2, R3, R4

의미

R1 = R2 * R3 + R4;

---

딥러닝에서 가장 많이 보이는 명령어 중 하나

왜냐면

sum += a*b;

가

FFMA

하나로 바뀐다.

---

GEMM 내부 대부분이

FFMA
FFMA
FFMA
FFMA

또는

HMMA
HMMA
HMMA

로 구성됨

---

3. Comparison 계열

---

FMAX

FMAX R1, R2, R3

의미

R1 = max(R2,R3);

---

ReLU

y=max(x,0);

↓

FMAX

---

FMNMX

Min/Max

최댓값
최솟값

선택

---

Reduction Max에서 자주 등장

---

4. Warp Communication

이게 GPU에서 매우 중요

---

SHFL

Shuffle

SHFL

---

다른 쓰레드 레지스터 읽기

예를 들어

x += __shfl_down_sync(mask,x,16);

↓

SHFL

---

Warp Reduction 핵심

sum
max
mean

거의 다 SHFL 등장

---

예시

for(offset=16;offset>0;offset/=2)
    x+=__shfl_down_sync(...);

↓

SHFL
FADD

SHFL
FADD

SHFL
FADD

---

이 패턴 보이면

Reduce Sum

가능성 높음

---

5. Special Function Unit

SFU

---

MUFU.EX2

MUFU.EX2

의미

2^x

---

exp(x)

계산 시

exp(x)
=
2^(x/log(2))

변환해서 사용

---

Softmax

Sigmoid

GELU

전부 등장

---

MUFU.RSQ

MUFU.RSQ

의미

1/sqrt(x)

---

LayerNorm

RMSNorm

BatchNorm

에서 거의 반드시 등장

---

패턴

sum
variance
RSQ
mul

보이면

Normalization

후보

---

6. Branch

---

BRA

Branch

BRA label

---

CPU의

goto

비슷

---

if문

loop

switch

전부 여기서 구현

---

ISETP

Set Predicate

ISETP

---

조건 계산

x<y

↓

ISETP

---

그 다음

@P0 BRA

형태

---

7. Tensor Core

딥러닝에서 제일 중요

---

HMMA

Hopper/Ampere MMA

HMMA

---

의미

D += A×B

---

Tensor Core 연산

---

만약

HMMA
HMMA
HMMA
HMMA
HMMA

수백줄

나온다

↓

거의 확실하게

GEMM
Attention
Conv
Linear

중 하나

---

실제로 SASS 읽을 때

처음엔 명령어 하나씩 보지 말고

이렇게 본다.

LDG
LDG
FFMA
FFMA
FFMA
STG

↓

데이터 읽음
곱셈누적
결과 저장

=
GEMM 후보

---

LDG
FMAX
SHFL
FMAX
SHFL

↓

최댓값 reduction

=
ReduceMax 후보

---

LDG
MUFU.EX2
SHFL
FADD
RCP
FMUL

↓

exp
sum
normalize

=
Softmax 후보