reduce_sum_f32 분석

좋습니다. reduce_sum_f32부터는 이전 LDG → primitive → STG 단순 elementwise 패턴과 달라졌습니다. 핵심 결론은 이겁니다.

reduce_sum_f32는 register-only 누적이 아니라,
shared memory에 중간값을 반복적으로 materialize하면서 reduction한다.

관찰된 핵심 명령어는:

STS       shared memory store
LDS       shared memory load
FADD      누적
BAR.SYNC  thread block synchronization
BRA       reduction loop
STG       final global store

입니다. 특히 초기값을 shared memory에 저장하고 BAR.SYNC로 동기화한 뒤 reduction loop에 들어가는 구조가 보입니다.

---

reduce_sum_f32 분석

문서 위치 추천:

notes/03_reduction/reduce_sum_f32.md

1. CUDA 코드 구조

우리가 작성한 커널은 대략 이런 구조입니다.

__global__ void reduce_sum_f32_kernel(const float* __restrict__ x,
                                      float* __restrict__ y,
                                      int n) {
    extern __shared__ float sdata[];

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

    float v = 0.0f;

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

    sdata[tid] = v;
    __syncthreads();

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sdata[tid] += sdata[tid + stride];
        }

        __syncthreads();
    }

    if (tid == 0) {
        y[blockIdx.x] = sdata[0];
    }
}

high-level 의미:

각 block이 x의 일부 구간을 shared memory로 가져온 뒤,
shared memory 안에서 tree reduction을 수행하고,
thread 0이 block sum을 y[blockIdx.x]에 저장한다.

---

2. 전체 SASS 역할 요약

SASS 흐름을 역할 단위로 압축하면 이렇게 됩니다.

R6 = blockIdx.x;
R7 = threadIdx.x;

R2 = blockIdx.x * blockDim.x + threadIdx.x;  // global i
R0 = 0.0f;

if (R2 < n) {
    R0 = x[R2];
}

sdata[tid] = R0;
barrier();

stride = blockDim.x >> 1;

while (stride != 0) {
    if (tid < stride) {
        lhs = sdata[tid];
        rhs = sdata[tid + stride];
        lhs = lhs + rhs;
        sdata[tid] = lhs;
    }

    barrier();
    stride >>= 1;
}

if (tid != 0) {
    exit;
}

R5 = sdata[0];
y[blockIdx.x] = R5;

이전 elementwise 커널과 결정적으로 다른 점은:

중간 결과가 register에만 머물지 않고 shared memory에 반복 저장된다.

입니다.

---

3. 초기 load + shared memory materialization

핵심 구간:

/*0010*/ S2R R6, SR_CTAID.X ;
/*0040*/ S2R R7, SR_TID.X ;
/*0050*/ IMAD R2, R6, c[0x0][0x0], R7 ;
/*0060*/ ISETP.GE.AND P0, PT, R2, c[0x0][0x170], PT ;
/*0070*/ @!P0 MOV R3, 0x4 ;
/*0080*/ @!P0 IMAD.WIDE R2, R2, R3, c[0x0][0x160] ;
/*0090*/ @!P0 LDG.E.CONSTANT R0, [R2.64] ;
/*00e0*/ STS [R7.X4], R0 ;
/*00f0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;

의미:

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

float v = 0.0f;

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

sdata[tid] = v;
__syncthreads();

여기서 R0이 중요합니다.

R0 = 0.0f
if (i < n) R0 = x[i]
sdata[tid] = R0

즉 범위 밖 thread는 0.0f를 shared memory에 넣습니다. 범위 안 thread만 global load를 수행합니다. 이 부분은 @!P0 LDG로 predicate 처리되어 있습니다.

---

4. STS [R7.X4], R0 의미

/*00e0*/ STS [R7.X4], R0 ;

여기서:

R7 = threadIdx.x
R7.X4 = threadIdx.x * 4
R0 = v

따라서:

sdata[tid] = v;

입니다.

이게 첫 번째 중요한 materialization boundary입니다.

global memory x[i]
→ register R0
→ shared memory sdata[tid]

즉 x[i]를 읽은 뒤 바로 누적하지 않고 shared memory에 저장합니다.

---

5. BAR.SYNC 의미

/*00f0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;

CUDA 코드의:

__syncthreads();

에 대응합니다.

reduction에서 이 barrier는 필수입니다.

모든 thread가 sdata[tid]를 채우기 전까지,
다음 단계의 sdata[tid + stride] load를 하면 안 되기 때문입니다.

SASS에서 반복적으로 보이는 BAR.SYNC.DEFER_BLOCKING이 바로 block-level synchronization 지점입니다. 초기 shared store 뒤와 loop 내부 store 뒤에 barrier가 나타납니다.

---

6. reduction loop 구조

핵심 loop 구간:

/*0130*/ ISETP.GE.AND P1, PT, R7, R3, PT ;
/*0140*/ @!P1 LEA R2, R3, R0, 0x2 ;
/*0150*/ @!P1 LDS R4, [R7.X4] ;
/*0160*/ SHF.R.U32.HI R3, RZ, 0x1, R3 ;
/*0170*/ @!P1 LDS R5, [R2] ;
/*0180*/ @!P1 FADD.FTZ R4, R4, R5 ;
/*0190*/ @!P1 STS [R7.X4], R4 ;
/*01a0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;
/*01b0*/ ISETP.NE.AND P1, PT, R3, RZ, PT ;
/*01c0*/ @P1 BRA 0x130 ;

이건 CUDA의 이 부분에 해당합니다.

for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
        sdata[tid] += sdata[tid + stride];
    }

    __syncthreads();
}

역할로 풀면:

if (tid < stride) {
    lhs = sdata[tid];
    rhs = sdata[tid + stride];
    lhs = lhs + rhs;
    sdata[tid] = lhs;
}

barrier();
stride >>= 1;

if (stride != 0) {
    goto loop;
}

LDS → LDS → FADD → STS가 reduction update의 핵심입니다.

---

7. shared memory accumulator update

reduction의 핵심 구간은 이것입니다.

/*0150*/ @!P1 LDS R4, [R7.X4] ;
/*0170*/ @!P1 LDS R5, [R2] ;
/*0180*/ @!P1 FADD.FTZ R4, R4, R5 ;
/*0190*/ @!P1 STS [R7.X4], R4 ;

의미:

float lhs = sdata[tid];
float rhs = sdata[tid + stride];
lhs = lhs + rhs;
sdata[tid] = lhs;

여기서 중요한 결론:

accumulator는 장기적으로 register에 유지되지 않는다.
각 stride마다 shared memory에서 읽고,
FADD 후 다시 shared memory에 저장된다.

즉 이 naive reduction은:

register accumulation

이라기보다:

shared memory materialized accumulation

입니다.

---

8. stride update

SASS에는 이런 명령이 보입니다.

/*00c0*/ USHF.R.U32.HI UR4, URZ, 0x1, UR4 ;
/*0160*/ SHF.R.U32.HI R3, RZ, 0x1, R3 ;

역할은 대략:

stride = blockDim.x >> 1;
stride >>= 1;

입니다.

처음에 blockDim.x를 읽고 절반으로 줄여 초기 stride를 만듭니다.

UR4 = blockDim.x >> 1

loop 안에서는:

R3 = R3 >> 1

로 다음 stride를 만듭니다.

USHF는 uniform register 쪽 shift, SHF는 일반 register 쪽 shift로 보면 됩니다.

---

9. tid < stride predicate

loop 안 조건은:

if (tid < stride)

입니다.

SASS에서는 반대 조건으로 predicate를 만듭니다.

/*0130*/ ISETP.GE.AND P1, PT, R7, R3, PT ;

의미:

P1 = (tid >= stride);

그 다음 실제 작업 명령어들은 모두 @!P1이 붙습니다.

@!P1 LDS ...
@!P1 LDS ...
@!P1 FADD ...
@!P1 STS ...

즉:

if (!(tid >= stride)) {
    // tid < stride
    sdata[tid] += sdata[tid + stride];
}

입니다.

여기서도 앞에서 봤던 패턴이 반복됩니다.

source 조건: tid < stride
SASS 조건: tid >= stride를 만든 뒤 @!P1로 실행

---

10. 최종 store

loop 종료 후:

/*01d0*/ @P0 EXIT ;
/*01e0*/ LDS R5, [RZ] ;
/*01f0*/ IMAD.MOV.U32 R3, RZ, RZ, 0x4 ;
/*0200*/ IMAD.WIDE.U32 R2, R6, R3, c[0x0][0x168] ;
/*0210*/ STG.E [R2.64], R5 ;
/*0220*/ EXIT ;

이 부분은 CUDA의:

if (tid == 0) {
    y[blockIdx.x] = sdata[0];
}

입니다.

앞에서:

/*00b0*/ ISETP.NE.AND P0, PT, R7, RZ, PT ;

로 P0 = (tid != 0)을 만들어둡니다.

그래서 loop가 끝난 뒤:

@P0 EXIT

즉:

if (tid != 0) return;

이 됩니다.

thread 0만 계속 진행해서:

R5 = sdata[0];
y[blockIdx.x] = R5;

를 수행합니다. 최종 결과는 shared memory에서 한 번 더 load된 뒤 global memory로 store됩니다.

---

11. register 흐름

R6

R6 = blockIdx.x
R6 = output block index

최종 store에서:

y[blockIdx.x]

주소 계산에 쓰입니다.

R7

R7 = threadIdx.x

대부분의 shared memory index로 쓰입니다.

STS [R7.X4], R0
LDS R4, [R7.X4]
STS [R7.X4], R4

즉:

sdata[tid]

입니다.

R2

초기: global index i
이후: x[i] address
loop 내부: sdata[tid + stride] address
최종: y[blockIdx.x] address

R2는 이 커널에서 의미가 가장 많이 바뀝니다.

R0

초기: 0.0f
조건부 load 후: x[i]
loop 내부: tid * 4 offset

초기에는 v 역할, 이후 loop에서는 주소 offset 계산에 재사용됩니다.

R3

초기: 4
loop: stride
최종: 4

R4

loop 초기: blockDim.x
loop 내부: sdata[tid] value
FADD 후: updated partial sum

R5

loop 내부: sdata[tid + stride]
최종: sdata[0]

---

12. 기존 elementwise와 비교

이전 커널들은 대략 이런 형태였습니다.

copy_global:
    LDG → STG

add_f32:
    LDG → LDG → FADD → STG

relu_f32:
    LDG → FMNMX → STG

clamp_f32:
    LDG → FSETP/FSEL → FSETP/FSEL → STG

하지만 reduce_sum_f32는 다릅니다.

reduce_sum_f32:
    LDG
    STS
    BAR.SYNC

    loop:
        LDS
        LDS
        FADD
        STS
        BAR.SYNC
        BRA

    LDS
    STG

즉 핵심 변화는:

global memory → register → shared memory → register → shared memory → ... → global memory

입니다.

---

13. materialization 관점 결론

이 커널은 우리가 세운 분석 프레임에서 매우 중요합니다.

CUDA source expression:
    sdata[tid] += sdata[tid + stride]

compiler graph rewrite:
    shared memory 기반 tree reduction loop

SASS primitive:
    LDS
    LDS
    FADD.FTZ
    STS
    BAR.SYNC
    BRA

intermediate materialization:
    있음.
    partial sum이 각 stride마다 shared memory에 저장됨.

특히 이 부분:

LDS R4, [R7.X4]
LDS R5, [R2]
FADD.FTZ R4, R4, R5
STS [R7.X4], R4

는 그대로:

sdata[tid] = sdata[tid] + sdata[tid + stride];

입니다.

그리고 STS가 있기 때문에 중간 partial sum은 register-only가 아니라 shared memory에 materialize됩니다.

---

14. 패턴 문서에 추가

notes/patterns/reduction_shared.md를 만들고 아래를 추가하면 됩니다.

# Pattern: Shared Memory Tree Reduction

## CUDA

```cpp
sdata[tid] = v;
__syncthreads();

for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
        sdata[tid] += sdata[tid + stride];
    }

    __syncthreads();
}

Observed SASS role pattern

STS [tid * 4], v
BAR.SYNC

loop:
    ISETP.GE predicate, tid, stride
    @active LDS lhs, [tid * 4]
    @active LDS rhs, [(tid + stride) * 4]
    @active FADD partial, lhs, rhs
    @active STS [tid * 4], partial
    BAR.SYNC
    stride >>= 1
    if stride != 0: BRA loop

Meaning

if (tid < stride) {
    sdata[tid] = sdata[tid] + sdata[tid + stride];
}

Observed in

• kernels/03_reduction/reduce_sum_f32.cu

Notes

• partial sum이 register에 장기 유지되지 않는다.
• 각 reduction step마다 shared memory에 저장된다.
• __syncthreads()는 BAR.SYNC.DEFER_BLOCKING으로 나타난다.
• 이 naive reduction은 shared memory materialization이 명확히 보이는 baseline이다.

---

## 15. 핵심 결론

`reduce_sum_f32`에서 얻은 결론은 이겁니다.

```text
단순 elementwise 커널은 대체로 LDG → primitive → STG 구조다.

하지만 naive reduction은:
LDG → STS → BAR.SYNC → LDS/LDS/FADD/STS 반복 → LDS → STG
구조로 내려간다.

그리고 가장 중요한 관찰:

partial sum은 register에 계속 누적되는 것이 아니라,
각 stride마다 shared memory에 저장된다.

이 커널은 앞으로 softmax_small_f32 분석의 기준선이 됩니다.
softmax의 max reduction, sum reduction도 결국 이와 비슷한 shared memory reduction 골격을 가질 가능성이 큽니다.