add_f32 분석

1. CUDA 코드

대상 커널 구조

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

    if (i < n) {
        float av = a[i];
        float bv = b[i];
        float out = av + bv;
        y[i] = out;
    }
}

SASS 의 역할 묶음으로 보면,

1. thread/block index 읽기
2. linear index i 계산
3. bounds check
4. a[i] 주소 계산
5. b[i] 주소 계산
6. a[i] load
7. b[i] load
8. FADD
9. y[i] 주소 계산
10. store

 

2. SASS 핵심 구조

SASS 의 핵심 부분

/*0000*/ MOV R1, c[0x0][0x28] ;
/*0010*/ S2R R6, SR_CTAID.X ;
/*0020*/ S2R R3, SR_TID.X ;
/*0030*/ IMAD R6, R6, c[0x0][0x0], R3 ;
/*0040*/ ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ;
/*0050*/ @P0 EXIT ;
/*0060*/ MOV R7, 0x4 ;
/*0070*/ ULDC.64 UR4, c[0x0][0x118] ;
/*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ;
/*0090*/ IMAD.WIDE R2, R6.reuse, R7.reuse, c[0x0][0x160] ;
/*00a0*/ LDG.E.CONSTANT R4, [R4.64] ;
/*00b0*/ LDG.E.CONSTANT R3, [R2.64] ;
/*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ;
/*00d0*/ FADD.FTZ R9, R4, R3 ;
/*00e0*/ STG.E [R6.64], R9 ;
/*00f0*/ EXIT ;

압축시 다음과 같이 변한다

R6 = blockIdx.x * blockDim.x + threadIdx.x
if R6 >= n: exit

R7 = 4

R4 = b_base + R6 * 4
R2 = a_base + R6 * 4

R4 = b[i]
R3 = a[i]

R6 = y_base + R6 * 4

R9 = R4 + R3

y[i] = R9

주의할 점으로 인자 순서,

CUDA 함수 인자가

const float* a,
const float* b,
float* y,
int n

라면 보통 constant memory argument 배치는 이런 식으로 해석할 수 있다.

c[0x0][0x160] = a base pointer
c[0x0][0x168] = b base pointer
c[0x0][0x170] = y base pointer
c[0x0][0x178] = n

따라서 이 부분은

IMAD.WIDE R4, R6, R7, c[0x0][0x168] ;
IMAD.WIDE R2, R6, R7, c[0x0][0x160] ;

각각

R4 = b + i * 4
R2 = a + i * 4

이다.

컴파일러가 a 주소보다 b 주소를 먼저 계산, CUDA 코드 순서와 반드시 같진 않음,

 

3. 명령어별 매핑

MOV R1, c[0x0][0x28]	런타임/스택 관련 초기화	직접 대응 없음
S2R R6, SR_CTAID.X	block index 읽기	blockIdx.x
S2R R3, SR_TID.X	thread index 읽기	threadIdx.x
IMAD R6, R6, c[0x0][0x0], R3	linear index 계산	i = blockIdx.x * blockDim.x + threadIdx.x
ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT	i >= n 비교	if (i < n)의 반대
@P0 EXIT	범위 밖 thread 종료	bounds check
MOV R7, 0x4	float byte size	sizeof(float)
IMAD.WIDE R4, R6, R7, c[0x0][0x168]	b[i] 주소 계산	&b[i]
IMAD.WIDE R2, R6, R7, c[0x0][0x160]	a[i] 주소 계산	&a[i]
LDG.E.CONSTANT R4, [R4.64]	b[i] load	float bv = b[i]
LDG.E.CONSTANT R3, [R2.64]	a[i] load	float av = a[i]
IMAD.WIDE R6, R6, R7, c[0x0][0x170]	y[i] 주소 계산	&y[i]
FADD.FTZ R9, R4, R3	float add	out = av + bv
STG.E [R6.64], R9	결과 저장	y[i] = out
EXIT	종료	함수 종료

 

4. copy_global 과 비교

copy 의 핵심 흐름,

IMAD.WIDE addr_x, i, 4, x_base
LDG val, [addr_x]
IMAD.WIDE addr_y, i, 4, y_base
STG [addr_y], val

add_f32 는 여기서 입력이 하나 더 늘어난 구조

IMAD.WIDE addr_b, i, 4, b_base
IMAD.WIDE addr_a, i, 4, a_base
LDG b_val, [addr_b]
LDG a_val, [addr_a]
IMAD.WIDE addr_y, i, 4, y_base
FADD out, b_val, a_val
STG [addr_y], out

즉 차이는

copy_global:
    LDG 1개
    STG 1개
    산술 연산 없음

add_f32:
    LDG 2개
    FADD 1개
    STG 1개

이것이 elementwise binary op 의 기본 패턴

 

5. index 계산 패턴

CUDA

int i = blockIdx.x * blockDim.x + threadIdx.x;

SASS

S2R R6, SR_CTAID.X ;
S2R R3, SR_TID.X ;
IMAD R6, R6, c[0x0][0x0], R3 ;

의미

R6 = blockIdx.x
R3 = threadIdx.x
R6 = R6 * blockDim.x + R3

결과적으로 

R6 = i

여기서 R6 이 이커널의 핵심 index register 이다.

 

6. .reuse 의미

새로 등장한 부분

IMAD.WIDE R2, R6.reuse, R7.reuse, c[0x0][0x160] ;

high-level 의미를 바꾸는 것은 아님.

R6.reuse = R6 값을 다시 사용한다
R7.reuse = R7 값을 다시 사용한다

즉 의미는 여전히

R2 = a_base + R6 * R7

하드웨어 스케줄링 / operand reuse cache 힌트에 가까움

같은 레지스터 값을 연속으로 사용할 때 레지스터 파일 접근 비용을 줄이기 위한 표시

 

레지스터를 고정된 변수명으로 보면 안 됨,

SASS 에서는 시점별 의미를 추적해야 한다. ( R4 는 처음에 주소였다가, 값 등으로 재사용 됨 )

 

7. FADD 분석

핵심 산술 명령어

FADD/FTZ R9, R4, R3 ;

 

의미

R9 = R4 + R3

이 시점에서

R4 = b[i]
R3 = a[i]

이므로

R9 = b[i] + a[i]

CUDA 코드로는

float out = av + bv;

FTZ 는 Flush To Zero

아주 작은 denormal/subnormal float 값을 0 으로 처리하는 모드

CMake 에서 --use_fast_math 와 같은 옵션을 사용하여 나옴

 

8. 이 분석에서 얻는 중요한 직관

add_f32는 앞으로 볼 대부분의 elementwise 연산의 기본형

예를 들어:

 

y[i] = a[i] * b[i];

 

이면 FADD만 FMUL로 바뀝니다.

LDG
LDG
FMUL
STG

 

그리고:

 

y[i] = a[i] * b[i] + c[i];

 

이면 보통:

LDG
LDG
LDG
FFMA
STG

 

또는:

LDG
LDG
LDG
FMUL
FADD
STG

 

가 됩니다.

그러니까 지금 정리한 add_f32 문서는 이후 mul_f32, fma_f32, bias_add, relu, layernorm 분석의 기준선이 됩니다.