copy_global 분석

1. CUDA 코드

대상 커널은 대략 이 구조이다.

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

    if (i < n) {
        y[i] = x[i];
    }
}

이 커널의 high-level 의미는 단순하다.

각 thread 가 하나의 index i 를 계산하고, i < n 이면 x[i] 를 읽어서 y[i] 에 저장한다.

즉 핵심 연산은 다음과 같다

y[i] = x[i];

SASS 에서는 이 한 줄이 보통 아래 역할들로 분해된다.

• thread / block index 읽기
• linear index i 계산
• bounds check
• x[i] 주소 계산
• x[i] load
• y[i] 주소 계산
• y[i] store
• exit

 

2. SASS 전체 구조 요약

SASS 핵심 부분은 다음이다.

/*0000*/ MOV R1, c[0x0][0x28] ;
/*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 ;
/*0070*/ ULDC.64 UR4, c[0x0][0x118] ;
/*0080*/ IMAD.WIDE R2, R4, R5, c[0x0][0x160] ;
/*0090*/ LDG.E.CONSTANT R3, [R2.64] ;
/*00a0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x168] ;
/*00b0*/ STG.E [R4.64], R3 ;
/*00c0*/ EXIT ;

역할만 압축하면

R4 = blockIdx.x * blockDim.x + threadIdx.x
P0 = (R4 >= n)
if P0 exit
R2 = x + R4 * 4
R3 = load x[i]
R4 = y + R4 * 4
store y[i] = R3
exit

역할만 압축하면

R4 = blockIdx.x * blockDim.x + threadIdx.x
P0 = (R4 >= n)
if P0 exit
R2 = x + R4 * 4
R3 = load x[i]
R4 = y + R4 * 4
store y[i] = R3
exit

 

3. 명령어별 매핑

• MOV R1, c[0x0][0x28]
  • 스택 포인터 / 런타임용 레지스터 초기화 - CUDA 코드 대응 없음
• S2R R4, SR_CTAID.x
  • blockIdx.x 읽기
    • blockIdx.x
• S2R R3, SR_TID.x
  • threadIdx.x 읽기
    • threadIdx.x
• IMAD R4, R4, c[0x0][0x0], R3
  • blockIdx.x * blockDim.x = threadIdx.x
    • int i - ...
• ISETP.GE.AND P0, PT, R4, x[0x0] [0x170], PT
  • i >= n 비교
    • if ( i < n ) 의 반대 조건
• @P0 EXIT
  • i >= n 이면 종료
    • bounds check
• MOV R5, 0x4
  • sizeof(float) = 4
    • float index byte offset 계산
• IMAX.WIDE R2, R4, R5, c[0x0][0x160]
  • x + i * 4 주소 계산
    • &y[i]
• STG.R [R4.64], R3
  • global memory 에 저장
    • y[i] = ...
• EXIT
  • 커널 종료
    • 함수 종료

 

4. CUDA 코드와 SASS 대응

CUDA

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

SASS

S2R R4, SR_CTAID.x ;
S2R R3, SR_TID.x ;
IMAD R4, R4, c[0x0][0x0], R3 ;

여기서 핵심은 IMAD 이다

R4 = R4 * c[0x0][0x0] + R3

각 값은 다음처럼 해석할 수 있다.

R4              = blockIdx.x
c[0x0][0x0]     = blockDim.x
R3              = threadIdx.x

그래서 결괒거으로

R4 = blockIdx.x * blockDim.x + threadIdx.x

즉, CUDA 의 i 가 SASS 에서는 R4 에 들어간다.

 

5. bounds check 분석

cuda

if (i < n) {
    y[i] = x[i];
}

SASS

ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT ;
@P0 EXIT ;

여기서 CUDA 코드는 i < n 인데 SASS 는 반대로 검사한다.

CUDA : if ( i < n ) 실행
SASS : if ( i >= n ) 종료

즉, 컴파일러가 조건을 이렇게 바꾼다.

if (i >= n) return;
y[i] = x[i];

이 형태가 SASS 에서는 더 자연스럽다.

P0 = (i >= n)
@P0 EXIT

P0 는 predicate register 이다. 일반 값 레지스터가 아니라 조건 실행을 위한 불리언 레지스터

 

6. 주소 계산 분석

CUDA

x[i]

SASS

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

여기서 R5 = 4 이다.

float 하나 = 4 bytes

따라서

R2 = x_base + i * 4

그 다음

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

는

float value = x[i];

에 대응된다.

R2.64 는 R2:R3 같은 64-bit 주소 표현으로 볼 수 있다. 주소는 64-bit 이고 , 데이터는 float 32-bit 이다. 

다만 이 SASS 에서는 LDG 결과가 R3 에 들어간다.

R3 = x[i]

 

7. store 분석

CUDA

y[i] = x[i];

SASS

IMAD.WIDE R4, R4, R5, c[0x0][0x168] ;
STG.E [R4.64], R3 ;

여기서는 y[i] 주소를 계산한다.

R4 = y_base + i * 4

그 다음

STG.E [R4.64], R3 ;

즉

global memory의 y[i] 위치에 R3 값을 저장

앞에서 R3 에는 x[i] 가 들어 있었다.

따라서 전체적으로

y[i] = x[i]

가 된다.

 

8. register 흐름

이 커널에서 핵심 register 흐름은 이렇게 볼 수 있다.

R4
 ├─ 처음: blockIdx.x
 ├─ IMAD 후: i = blockIdx.x * blockDim.x + threadIdx.x
 ├─ bounds check에 사용
 ├─ x 주소 계산에 사용
 └─ 나중에 y 주소 계산 결과로 덮어쓰기됨
 
 R3
 ├─ 처음: threadIdx.x
 └─ LDG 후: x[i] 값으로 덮어쓰기됨
 
 R5
 └─ 4, 즉 sizeof(float)
 
 R2
 └─ x[i]의 64-bit global address

중요한 관찰은 이것이다.

• 같은 register 가 계속 같은 의미를 유지하지 않는다.

그래서 SASS 분석에서는 R3는 무엇이다라고 고정하면 안되다.

정확히는 이 시점의 R3 는 무엇이다라고 봐야 한다.

 

9. high-level pseudo SASS

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

R4 = R4 * blockDim.x + R3;

P0 = R4 >= n;
if (P0) exit;

R5 = 4;

R2 = x_base + R4 * R5;
R3 = global_load_float(R2);

R4 = y_base + R4 * R5;
global_store_float(R4, R3);

exit;

 

10. 여기서 얻는 기본 패턴

copy_global 에서 얻는 첫 번째 SASS 패턴은 이것이다.

Pattern: 1D elementwise kernel index 계산

S2R R?, SR_CTAID.X
S2R R?, SR_TID.X
IMAD R_i, R_cta, blockDim.x, R_tid

의미

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

두 번째 패턴

Pattern: bounds check

ISETP.GE.AND P0, PT, R_i, n, PT
@P0 EXIT

의미

if (i >= n) return;

CUDA 원본과 반대 조건으로 exit 하는 조건이 SASS 에서 자주 나온다..

세 번째 패턴

Pattern: float global load

MOV R_stride, 0x4
IMAD.WIDE R_addr, R_i, R_stride, x_base
LDG.E R_val, [R_addr.64]

의미

float val = x[i];

네 번째 패턴

Pattern: float global store

IMAD.WIDE R_addr, R_i, R_stride, y_base
STG.E [R_addr.64], R_val

의미

y[i] = val;

 

11. copy_global 에서 특히 중요한 관찰

• special register 읽기
• linear thread index 계산
• predicate 기반 bounds check
• byte offset 계산
• global load
• global store
• register  reuse