mul_f32 분석

1. CUDA 코드

대상 커널 구조

__global__ void mul_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;
    }
}

high-level 의미

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

SASS 역할 묶음

1. blockIdx.x / threadIdx.x 읽기
2. i = blockIdx.x * blockDim.x + threadIdx.x
3. i >= n이면 종료
4. b[i] 주소 계산
5. a[i] 주소 계산
6. b[i] load
7. a[i] load
8. y[i] 주소 계산
9. FMUL
10. y[i] store

 

2. SASS 핵심 구조

/*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 ;
/*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*/ FMUL.FTZ R9, R4, R3 ;
/*00e0*/ STG.E [R6.64], R9 ;
/*00f0*/ EXIT ;

이는

R6 = blockIdx.x;
R3 = threadIdx.x;

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

if (R6 >= n) {
    exit;
}

R7 = 4;                    // sizeof(float)

R4 = b_base + R6 * 4;       // &b[i]
R2 = a_base + R6 * 4;       // &a[i]

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

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

R9 = R4 * R3;               // b[i] * a[i]

store y[i] = R9;

 

3. 명령어별 매핑

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 stride	sizeof(float)
IMAD.WIDE R4, R6, R7, c[0x0][0x168]	b[i] 주소 계산	&b[i]
IMAD.WIDE R2, R6.reuse, R7.reuse, 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]
FMUL.FTZ R9, R4, R3	float multiply	out = av * bv
STG.E [R6.64], R9	결과 저장	y[i] = out

 

4. add_f32 와의 직접 비교

add_f32 와 mul_f32 는 거의 동일

add_f32

LDG.E.CONSTANT R4, [R4.64] ;
LDG.E.CONSTANT R3, [R2.64] ;
IMAD.WIDE R6, R6, R7, c[0x0][0x170] ;
FADD.FTZ R9, R4, R3 ;
STG.E [R6.64], R9 ;

mul_f32

LDG.E.CONSTANT R4, [R4.64] ;
LDG.E.CONSTANT R3, [R2.64] ;
IMAD.WIDE R6, R6, R7, c[0x0][0x170] ;
FMUL.FTZ R9, R4, R3 ;
STG.E [R6.64], R9 ;

차이는 딱 하나

• add_f32 : FADD.FTZ
• mul_f32 : FMUl.FTZ

즉 SASS 관점에서 보면

y[i] = a[i] + b[i]

는

LDG, LDG, FADD, STG

이고

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

는

LDG, LDG, FMUl, STG

이다.

이러한 기준점은

나중에 더 복잡한 operator 를 볼 때, 이런 산술 명령어 하나가 어떤 high-level 역학을 가졌는지 판단하는 기본 단위가 된다.

 

5. 레지스터 호출

R6

• R6 = blockIdx.x
• R6 = i
• R6 = y[i] address

처음엔 block index, 그 다음엔 logical index, 마지막엔 output address 이다.

R3

• R3 = threadIdx.x
• R3 = a[i]

R3 은 thread index 로 쓰였다가 load 결과값으로 덮인다

...

등등 각 레지스터 존재