fma_f32 분석

1. CUDA 코드

대상 커널 구조

__global__ void fma_f32_kernel(const float* __restrict__ a,
                               const float* __restrict__ b,
                               const float* __restrict__ c,
                               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 cv = c[i];
        float out = av * bv + cv;
        y[i] = out;
    }
}

SASS 역할 묶음

• blockIdx.x / threadIdx.x 읽기
• i = blockIdx.x * blockDim.x + threadIdx.x
• i >= n 이면 종료
• b[i] 주소 계산
• a[i] 주소 계산
• b[i] load
• c[i] 주소 계산
• a[i] load
• c[i] load
• y[i] 주소 계산
• FFMA
• y[i] store

 

2. SASS 핵심 구조

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

/*0060*/ MOV R9, 0x4 ;

/*0080*/ IMAD.WIDE R4, R8, R9, c[0x0][0x168] ;
/*0090*/ IMAD.WIDE R2, R8.reuse, R9.reuse, c[0x0][0x160] ;
/*00a0*/ LDG.E.CONSTANT R4, [R4.64] ;

/*00b0*/ IMAD.WIDE R6, R8.reuse, R9.reuse, c[0x0][0x170] ;
/*00c0*/ LDG.E.CONSTANT R3, [R2.64] ;
/*00d0*/ LDG.E.CONSTANT R7, [R6.64] ;

/*00e0*/ IMAD.WIDE R8, R8, R9, c[0x0][0x178] ;
/*00f0*/ FFMA.FTZ R11, R4, R3, R7 ;
/*0100*/ STG.E [R8.64], R11 ;
/*0110*/ EXIT ;


R8 = blockIdx.x;
R3 = threadIdx.x;

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

if (R8 >= n) {
    exit;
}

R9 = 4;                      // sizeof(float)

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

R4 = load b[i];

R6 = c_base + R8 * 4;         // &c[i]

R3 = load a[i];
R7 = load c[i];

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

R11 = R4 * R3 + R7;           // b[i] * a[i] + c[i]

store y[i] = R11;

 

3. 인자 배치 추정

CUDA 함수 시그니처가

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

이것의 constant memory argument 배치

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

실제 SASS 흐름

IMAD.WIDE R4, R8, R9, c[0x0][0x168] ;  // b[i] address
IMAD.WIDE R2, R8, R9, c[0x0][0x160] ;  // a[i] address
IMAD.WIDE R6, R8, R9, c[0x0][0x170] ;  // c[i] address
IMAD.WIDE R8, R8, R9, c[0x0][0x178] ;  // y[i] address

CUDA 코드상으로는 a, b, c, y 순서이지만 SASS 에서는 주소 계산과 load 순서가 약간 바뀌었다. 

컴파일러가 dependency 와 schduiling 을 보고 재배치한다.

 

4. add / mul / fma 직접 비교

add_f32

FADD.FTZ R9, R4, R3 ;

out = b[i] + a[i];

 

mul_f32

FMUL.FTZ R9, R4, R3 ;

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

 

fma_f32

FFMA.FTZ R11, R4, R3, R7 ;

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

 

정리하면

CUDA source              SASS core instruction
------------------------------------------------
a + b                    FADD.FTZ
a * b                    FMUL.FTZ
a * b + c                FFMA.FTZ

 

2개의 연산처럼 보이는 것이 FFMA 로 내려간 것을 확인

즉 SASS 는 source-level syntax 를 그대로 보존하지 않는다.

컴파일러가 수학적으로 허용되는 범위에서 더 낮은 primitive로 재구성한다.