좋습니다. clamp_f32는 예상과 다르게 FMNMX 두 개가 아니라 FSETP + FSEL 두 쌍으로 내려갔습니다.
핵심 구간은 여기입니다.
/*00b0*/ FSETP.GEU.FTZ.AND P0, PT, R2, c[0x0][0x170], PT ;
/*00c0*/ FSEL R0, R2, c[0x0][0x170], P0 ;
/*00d0*/ FSETP.GT.FTZ.AND P0, PT, R0, c[0x0][0x174], PT ;
/*00e0*/ FSEL R7, R0, c[0x0][0x174], !P0 ;
즉 source-level의:
if (v < lo) {
v = lo;
}
if (v > hi) {
v = hi;
}
가 SASS에서는 branch 없이:
tmp = (v >= lo) ? v : lo;
out = (tmp > hi) ? hi : tmp;
형태로 내려갔습니다.
clamp_f32 분석
문서 위치 추천:
notes/02_control/clamp_f32.md
1. CUDA 코드
대상 커널은 대략 이 구조입니다.
__global__ void clamp_f32_kernel(const float* __restrict__ x,
float* __restrict__ y,
float lo,
float hi,
int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float v = x[i];
if (v < lo) {
v = lo;
}
if (v > hi) {
v = hi;
}
y[i] = v;
}
}
high-level 의미는 이것입니다.
y[i] = min(max(x[i], lo), hi);
2. SASS 핵심 구조
/*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][0x178], PT ;
/*0050*/ @P0 EXIT ;
/*0060*/ IMAD.MOV.U32 R5, RZ, RZ, 0x4 ;
/*0080*/ IMAD.WIDE R2, R4, R5, c[0x0][0x160] ;
/*0090*/ LDG.E.CONSTANT R2, [R2.64] ;
/*00a0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x168] ;
/*00b0*/ FSETP.GEU.FTZ.AND P0, PT, R2, c[0x0][0x170], PT ;
/*00c0*/ FSEL R0, R2, c[0x0][0x170], P0 ;
/*00d0*/ FSETP.GT.FTZ.AND P0, PT, R0, c[0x0][0x174], PT ;
/*00e0*/ FSEL R7, R0, c[0x0][0x174], !P0 ;
/*00f0*/ STG.E [R4.64], R7 ;
/*0100*/ EXIT ;
사람이 읽기 쉽게 바꾸면:
R4 = blockIdx.x;
R3 = threadIdx.x;
R4 = R4 * blockDim.x + R3; // i
if (R4 >= n) {
exit;
}
R5 = 4; // sizeof(float)
R2 = x_base + R4 * 4; // &x[i]
R2 = load x[i]; // v = x[i]
R4 = y_base + R4 * 4; // &y[i]
// lower clamp
P0 = R2 >= lo;
R0 = P0 ? R2 : lo;
// upper clamp
P0 = R0 > hi;
R7 = !P0 ? R0 : hi;
store y[i] = R7;
즉 최종적으로:
R7 = min(max(x[i], lo), hi);
입니다.
3. 인자 배치 추정
CUDA 함수 시그니처가:
const float* x,
float* y,
float lo,
float hi,
int n
이므로 constant memory argument는 이렇게 해석할 수 있습니다.
c[0x0][0x160] = x base pointer
c[0x0][0x168] = y base pointer
c[0x0][0x170] = lo
c[0x0][0x174] = hi
c[0x0][0x178] = n
SASS와도 맞습니다.
IMAD.WIDE R2, R4, R5, c[0x0][0x160] ; // x[i] address
IMAD.WIDE R4, R4, R5, c[0x0][0x168] ; // y[i] address
FSETP.GEU.FTZ.AND P0, PT, R2, c[0x0][0x170], PT ; // v >= lo
FSETP.GT.FTZ.AND P0, PT, R0, c[0x0][0x174], PT ; // tmp > hi
4. 명령어별 매핑
| S2R R4, SR_CTAID.X | block index 읽기 | blockIdx.x |
| S2R R3, SR_TID.X | thread index 읽기 | threadIdx.x |
| IMAD R4, R4, c[0x0][0x0], R3 | linear index 계산 | i = blockIdx.x * blockDim.x + threadIdx.x |
| ISETP.GE.AND P0, PT, R4, c[0x0][0x178], PT | i >= n 비교 | bounds check |
| @P0 EXIT | 범위 밖 thread 종료 | if (i >= n) return |
| IMAD.MOV.U32 R5, RZ, RZ, 0x4 | float byte stride | sizeof(float) |
| IMAD.WIDE R2, R4, R5, c[0x0][0x160] | x[i] 주소 계산 | &x[i] |
| LDG.E.CONSTANT R2, [R2.64] | x[i] load | float v = x[i] |
| IMAD.WIDE R4, R4, R5, c[0x0][0x168] | y[i] 주소 계산 | &y[i] |
| FSETP.GEU.FTZ.AND P0, PT, R2, lo, PT | v >= lo 비교 | lower clamp 조건 |
| FSEL R0, R2, lo, P0 | P0 ? v : lo | max(v, lo) |
| FSETP.GT.FTZ.AND P0, PT, R0, hi, PT | tmp > hi 비교 | upper clamp 조건 |
| FSEL R7, R0, hi, !P0 | !P0 ? tmp : hi | min(tmp, hi) |
| STG.E [R4.64], R7 | 결과 저장 | y[i] = clamp(v, lo, hi) |
SASS의미CUDA 대응
5. lower clamp 분석
CUDA 원본:
if (v < lo) {
v = lo;
}
수학적으로는:
v = max(v, lo);
SASS:
FSETP.GEU.FTZ.AND P0, PT, R2, c[0x0][0x170], PT ;
FSEL R0, R2, c[0x0][0x170], P0 ;
여기서:
R2 = v = x[i]
c[0x0][0x170] = lo
첫 줄:
P0 = (v >= lo);
두 번째 줄:
R0 = P0 ? v : lo;
즉:
R0 = max(v, lo);
입니다.
중요한 점은 source에는 if (v < lo)였는데, SASS는 반대 조건인 v >= lo로 선택합니다.
source:
if (v < lo) v = lo
SASS:
if (v >= lo) select v
else select lo
결과는 같지만 표현은 다릅니다.
6. upper clamp 분석
CUDA 원본:
if (v > hi) {
v = hi;
}
수학적으로는:
v = min(v, hi);
SASS:
FSETP.GT.FTZ.AND P0, PT, R0, c[0x0][0x174], PT ;
FSEL R7, R0, c[0x0][0x174], !P0 ;
여기서:
R0 = max(x[i], lo)
c[0x0][0x174] = hi
첫 줄:
P0 = (R0 > hi);
두 번째 줄은 predicate가 !P0입니다.
R7 = !P0 ? R0 : hi;
즉:
if (R0 <= hi) {
R7 = R0;
} else {
R7 = hi;
}
결과적으로:
R7 = min(R0, hi);
입니다.
따라서 전체는:
R7 = min(max(x[i], lo), hi);
7. ReLU와 clamp 비교
relu_f32에서는:
FMNMX.FTZ R7, RZ, R2, !PT ;
단일 min/max instruction이 나왔습니다.
반면 clamp_f32에서는:
FSETP.GEU.FTZ.AND P0, PT, R2, lo, PT ;
FSEL R0, R2, lo, P0 ;
FSETP.GT.FTZ.AND P0, PT, R0, hi, PT ;
FSEL R7, R0, hi, !P0 ;
즉 clamp는 FMNMX 두 개가 아니라:
compare + select
compare + select
로 내려갔습니다.
정리하면:
relu_f32:
max(x, 0)
→ FMNMX
clamp_f32:
min(max(x, lo), hi)
→ FSETP + FSEL + FSETP + FSEL
이 차이가 중요합니다.
왜 ReLU는 FMNMX인데 clamp는 FMNMX 두 개가 아니냐는 질문이 생깁니다. 가능한 이유는 lo, hi가 kernel argument로 들어온 값이고, source가 명시적인 if 두 개였기 때문입니다. 컴파일러가 이 경우에는 min/max instruction으로 축약하지 않고 predicate select 형태를 택했습니다.
8. branch가 없는 점이 핵심
clamp_f32는 source-level에서는 명확한 branch 코드입니다.
if (v < lo) {
v = lo;
}
if (v > hi) {
v = hi;
}
하지만 SASS에는 내부 분기 BRA가 없습니다.
실제 조건 처리 부분은:
FSETP
FSEL
FSETP
FSEL
입니다.
즉 warp divergence를 유발하는 branch가 아니라, predicate 기반 select로 처리되었습니다.
이 결론이 중요합니다.
source-level if가 항상 SASS-level branch가 되는 것은 아니다.
짧은 값 선택 로직은 branchless select로 lowering될 수 있다.
9. FSETP와 FSEL 의미
FSETP
FSETP.GEU.FTZ.AND P0, PT, R2, lo, PT ;
floating-point compare 후 predicate register를 설정합니다.
P0 = (R2 >= lo);
GEU는 greater-or-equal unordered 계열 비교로 보면 됩니다. 여기서는 실질적으로 >= 비교입니다.
FSEL
FSEL R0, R2, lo, P0 ;
predicate에 따라 두 float 값 중 하나를 선택합니다.
R0 = P0 ? R2 : lo;
즉 FSETP + FSEL은 high-level의 짧은 조건식에 대응됩니다.
out = cond ? a : b;
10. 레지스터 흐름
R4
├─ blockIdx.x
├─ i
└─ y[i] address
R3
└─ threadIdx.x
R2
├─ x[i] address
└─ x[i] value
R5
└─ 4, sizeof(float)
R0
└─ max(x[i], lo)
R7
└─ min(max(x[i], lo), hi)
핵심 흐름:
R2 = x[i]
R0 = max(R2, lo)
R7 = min(R0, hi)
STG y[i] = R7
11. 패턴 문서에 추가
notes/patterns/control_f32.md에 이어서 추가하면 됩니다.
## Pattern: F32 Clamp
### CUDA
```cpp
if (v < lo) {
v = lo;
}
if (v > hi) {
v = hi;
}
Equivalent:
v = min(max(v, lo), hi);
Observed SASS
FSETP.GEU.FTZ.AND P0, PT, R_v, R_lo, PT ;
FSEL R_tmp, R_v, R_lo, P0 ;
FSETP.GT.FTZ.AND P0, PT, R_tmp, R_hi, PT ;
FSEL R_out, R_tmp, R_hi, !P0 ;
Meaning
R_tmp = (R_v >= lo) ? R_v : lo;
R_out = (R_tmp > hi) ? hi : R_tmp;
Equivalent:
R_out = min(max(R_v, lo), hi);
Observed in
- kernels/02_control/clamp_f32.cu
Notes
- source-level if 두 개가 SASS-level branch로 내려가지 않았다.
- 실제 SASS에는 FSETP + FSEL 두 쌍이 나타났다.
- ReLU는 FMNMX 하나였지만, clamp는 FSETP/FSEL 기반 branchless select로 표현되었다.
- 이 패턴은 thresholding, clipping, activation bound, quantization clamp 등에서 반복적으로 나타날 수 있다.
---
## 12. 현재까지 control 계열 결론
이제 control 계열에서 중요한 관찰 두 개가 생겼습니다.
```text
relu_f32:
v > 0 ? v : 0
→ FMNMX.FTZ
clamp_f32:
if (v < lo) v = lo;
if (v > hi) v = hi;
→ FSETP + FSEL + FSETP + FSEL
즉 둘 다 source-level 조건을 포함하지만, SASS에서는 실제 branch가 아닙니다.
조건문처럼 보이는 CUDA 코드
→ compiler graph rewrite
→ branchless SASS primitive
이 프레임에 넣으면:
CUDA source expression:
if (v < lo) v = lo;
if (v > hi) v = hi;
compiler graph rewrite:
tmp = max(v, lo)
out = min(tmp, hi)
SASS primitive:
FSETP + FSEL
FSETP + FSEL
intermediate materialization:
global memory materialization 없음
R0 register에 lower-clamped intermediate가 잠깐 존재
최종 R7이 store됨
여기서 “intermediate materialization” 관점도 좋습니다.
R0에 max(x, lo)라는 중간값이 생기지만, 이건 register 안에만 있습니다.
R0 = max(x[i], lo)
R7 = min(R0, hi)
STG y[i] = R7
즉 lower clamp 결과는 global/shared memory에 저장되지 않습니다.
13. 다음 단계
이제 control까지 충분히 봤습니다. 다음은 reduce_sum_f32.sass가 맞습니다.
거기서부터는 관점이 바뀝니다.
지금까지는:
LDG → primitive → STG
였는데, reduction은:
LDG
→ STS shared memory
→ BAR.SYNC
→ LDS
→ FADD
→ STS
→ 반복
→ STG
형태가 됩니다.
봐야 할 핵심은 이겁니다.
1. accumulator가 register에 머무는 구간
2. shared memory에 materialize되는 구간
3. __syncthreads()가 BAR.SYNC로 나타나는 위치
4. 최종 global store boundary
reduce_sum_f32.sass를 붙여주면 다음엔 memory hierarchy 관점으로 분석하면 됩니다.
'SASS_Probe' 카테고리의 다른 글
| softmax_small_f32 분석 (0) | 2026.06.14 |
|---|---|
| reduce_sum_f32 분석 (0) | 2026.06.14 |
| relu_f32 분석 (0) | 2026.06.14 |
| fma_contract_f32: FFMA Contraction Pattern Experiment (0) | 2026.06.13 |
| fma_f32 분석 (0) | 2026.06.13 |