기본 연산 의미 명세 07 - Maximum과 Minimum

1. 연산 개요

Maximum과 Minimum은 두 입력값을 비교하여 더 큰 값 또는 더 작은 값을 출력하는 연산이다.

기본 형태는 다음과 같다.

maximum:
y[i] = max(x[i], z[i])

minimum:
y[i] = min(x[i], z[i])

Maximum의 의미는 다음과 같다.

x[i] >= z[i]이면 y[i] = x[i]
그렇지 않으면 y[i] = z[i]

Minimum은 반대다.

x[i] <= z[i]이면 y[i] = x[i]
그렇지 않으면 y[i] = z[i]

데이터 의존성은 다음과 같다.

x[i] ─┐
      ├→ Compare and Select → y[i]
z[i] ─┘

수학적으로 Maximum과 Minimum은 Comparison과 Selection의 합성으로 해석할 수 있다.

max(x, z)
=
select(x >= z, x, z)

min(x, z)
=
select(x <= z, x, z)

하지만 실제 GPU에서는 Comparison과 Selection을 별도로 수행하지 않고, 직접 maximum 또는 minimum 명령으로 lowering될 수 있다.

Maximum과 Minimum은 다음 연산의 기초가 된다.

ReLU
Clamp
Max pooling
Min pooling
Max reduction
Min reduction
Saturation
Thresholding
Bounding box 처리
Numerical stabilization
Softmax의 최대값 계산
Argmax의 값 비교
Gradient clipping
Quantization 범위 제한

2. 기본 수학적 정의

두 실수 x, z에 대해 Maximum은 다음과 같이 정의한다.

max(x, z)
=
x, if x >= z
z, otherwise

Minimum은 다음과 같다.

min(x, z)
=
x, if x <= z
z, otherwise

Elementwise tensor 형태에서는 다음과 같다.

y[i] = max(x[i], z[i])

w[i] = min(x[i], z[i])

동일한 shape의 두 입력을 사용한다면:

shape(X) = shape(Z) = shape(Y)

예:

X = [1, 5, -2]
Z = [3, 2, -4]

Maximum:

max(X, Z) = [3, 5, -2]

Minimum:

min(X, Z) = [1, 2, -4]

3. Scalar Maximum과 Minimum

입력 tensor의 각 원소를 하나의 scalar와 비교할 수 있다.

y[i] = max(x[i], c)

y[i] = min(x[i], c)

대표적인 예는 ReLU다.

y[i] = max(x[i], 0)

Upper bound 제한은 다음과 같다.

y[i] = min(x[i], upper)

Lower bound 제한은 다음과 같다.

y[i] = max(x[i], lower)

Scalar c는 모든 출력 원소에서 재사용된다.

c → y[0]
  → y[1]
  → y[2]
  → ...

실행에서는 scalar가 다음 방식으로 공급될 수 있다.

Immediate constant
Kernel argument
Register
Constant memory
Uniform register

4. Broadcasting Maximum과 Minimum

입력 shape이 다르면 broadcasting을 적용할 수 있다.

예를 들어 channel별 lower bound를 적용할 수 있다.

y[n, c] = max(x[n, c], lower[c])

또는 channel별 upper bound:

y[n, c] = min(x[n, c], upper[c])

lower[c]와 upper[c]는 batch 축 전체에서 반복 사용된다.

lower[c]
→ y[0, c]
→ y[1, c]
→ y[2, c]

Broadcasting에서는 다음 정보가 연산 의미에 포함된다.

Bound가 적용되는 축
Broadcast index mapping
Bound의 dtype
Bound의 memory layout
같은 값을 공유하는 thread 범위

5. Elementwise Maximum과 Max Reduction의 구분

Maximum이라는 동일한 관계를 사용해도 elementwise maximum과 max reduction은 다른 연산이다.

Elementwise Maximum

y[i] = max(x[i], z[i])

구조:

두 대응 원소
→ 출력 원소 하나

출력 원소 사이에 의존성이 없다.

Max Reduction

m = max(x[0], x[1], ..., x[N-1])

구조:

여러 입력 원소
→ 집계된 최대값 하나

Max reduction은 cross-element dependency를 가진다.

실행 구조도 다르다.

Elementwise Maximum:

LDG X
LDG Z
FMAX
STG Y

Max Reduction:

여러 입력 load
→ thread-local maximum
→ warp/block reduction
→ 최종 maximum store

따라서 SASS에서 FMAX 하나를 발견했다고 해서 elementwise maximum인지 max reduction인지 판단할 수 없다.

반복 구조, shuffle, shared memory, accumulator 재사용을 함께 봐야 한다.

6. 데이터 의존성

Elementwise Maximum은 다음과 같다.

y[i] = max(x[i], z[i])

출력 y[i]는 같은 위치의 두 입력에 의존한다.

x[i] ─┐
      ├→ y[i]
z[i] ─┘

다른 위치의 입력에는 의존하지 않는다.

i != j일 때
y[i]는 x[j], z[j]를 필요로 하지 않는다.

따라서 각 출력은 독립적으로 병렬 실행할 수 있다.

thread 0 → y[0]
thread 1 → y[1]
thread 2 → y[2]

Max reduction에서는 이 독립성이 사라진다.

m = max_i(x[i])

모든 입력이 하나의 결과에 영향을 준다.

7. 의미 불변성

Maximum과 Minimum의 구현이 달라져도 다음 조건은 유지되어야 한다.

7.1 순서 관계

Maximum은 정의된 비교 의미에 따라 더 큰 값을 선택해야 한다.

x > z이면 max(x, z) = x
x < z이면 max(x, z) = z

Minimum은 더 작은 값을 선택한다.

7.2 동일값 처리

두 입력이 수치적으로 같을 때 어느 operand를 반환하는지 구분해야 할 수 있다.

x == z

일반 실수에서는 어느 값을 선택해도 결과값은 같다.

하지만 floating-point에서는 다음이 다를 수 있다.

+0.0과 -0.0
서로 다른 NaN payload
동일 수치값이지만 다른 bit pattern

따라서 엄격한 의미에서는 tie-breaking 규칙도 명세할 필요가 있다.

7.3 위치 대응 관계

Elementwise 연산에서는 같은 논리적 위치의 두 값을 비교해야 한다.

y[i] = max(x[i], z[i])

7.4 Broadcasting axis

Broadcasting이 있다면 어느 축의 bound를 사용하는지 보존해야 한다.

7.5 특수값 의미

NaN과 signed zero 처리 방식은 연산 의미의 일부다.

8. 대수적 성질

8.1 교환법칙

일반 실수에서는 다음이 성립한다.

max(x, z) = max(z, x)

min(x, z) = min(z, x)

하지만 NaN payload 선택이나 signed zero tie-breaking까지 bitwise하게 고려하면 operand 순서가 결과 representation에 영향을 줄 가능성이 있다.

8.2 결합법칙

일반 실수에서는 다음이 성립한다.

max(max(x, z), w)
=
max(x, max(z, w))

min(min(x, z), w)
=
min(x, min(z, w))

이 결합성은 max/min reduction을 tree 형태로 병렬화하는 근거가 된다.

sequential:
max(max(max(x0, x1), x2), x3)

tree:
max(max(x0, x1), max(x2, x3))

하지만 NaN 처리 규칙과 tie-breaking 규칙에 따라 bitwise 결과는 달라질 수 있다.

8.3 멱등성

같은 값을 두 번 적용하면 결과는 그대로다.

max(x, x) = x

min(x, x) = x

이 성질은 중복 연산 제거에 사용할 수 있다.

다만 x = NaN이고 연산이 특정 NaN canonicalization을 수행한다면 bitwise 결과를 별도로 확인해야 한다.

8.4 흡수 법칙

Maximum과 Minimum은 다음 관계를 가진다.

max(x, min(x, z)) = x

min(x, max(x, z)) = x

이 성질은 중첩된 bound 연산을 단순화할 수 있는 근거가 된다.

8.5 순서 단조성

x <= z이면 다음이 성립한다.

max(x, c) <= max(z, c)

min(x, c) <= min(z, c)

즉 Maximum과 Minimum은 각 operand에 대해 단조적이다.

이 성질은 범위 분석과 안정적인 bound propagation에 사용할 수 있다.

9. Bounding 의미

Scalar Maximum은 lower bound를 적용한다.

y = max(x, lower)

결과는 반드시 다음을 만족한다.

y >= lower

Scalar Minimum은 upper bound를 적용한다.

y = min(x, upper)

결과는 다음을 만족한다.

y <= upper

두 연산을 결합하면 Clamp가 된다.

y = min(max(x, lower), upper)

결과 범위:

lower <= y <= upper

단, lower <= upper라는 조건이 필요하다.

10. NaN 처리 의미

Floating-point Maximum과 Minimum에서 가장 중요한 문제는 NaN이다.

NaN은 일반적인 순서 관계에 참여하지 않는다.

NaN < x  → false
NaN > x  → false
NaN == x → false

따라서 Maximum과 Minimum에는 서로 다른 두 의미가 존재할 수 있다.

10.1 NaN 전파형

입력 중 하나가 NaN이면 결과도 NaN이다.

max_propagate(NaN, x) = NaN
min_propagate(NaN, x) = NaN

10.2 숫자 우선형

한 입력만 NaN이면 다른 숫자 operand를 반환한다.

max_number(NaN, x) = x
min_number(NaN, x) = x

두 입력이 모두 NaN이면 NaN을 반환한다.

이 두 의미는 일반적인 유한 입력에서는 같지만 NaN에서 다르다.

따라서 다음 표현들이 반드시 같은 것은 아니다.

max(x, z)

select(x >= z, x, z)

예를 들어 x = NaN, z = 1이면 comparison은 false다.

x >= z → false

따라서 selection은 z를 반환한다.

select(false, NaN, 1) = 1

이는 숫자 우선형 maximum과는 맞지만 NaN 전파형 maximum과는 다르다.

11. NaN과 Operand 순서

다음 표현을 생각하자.

select(x >= z, x, z)

x = 1, z = NaN이면:

1 >= NaN → false

따라서 NaN이 반환된다.

select(false, 1, NaN) = NaN

반대로:

select(z <= x, x, z)

에서도 ordered comparison의 세부 규칙에 따라 같은 결과가 보장되지 않을 수 있다.

즉 comparison-selection으로 max/min을 재작성할 때는 다음을 명확히 해야 한다.

NaN 전파 여부
Ordered/unordered comparison
Operand 순서
Tie-breaking
Signed zero 처리

12. Signed Zero

Floating-point에는 +0.0과 -0.0이 존재한다.

일반적인 comparison에서는 두 값이 같다.

+0.0 == -0.0 → true

하지만 후속 연산에서는 차이가 있다.

1 / +0.0 = +Inf
1 / -0.0 = -Inf

Maximum과 Minimum의 바람직한 결과는 일반적으로 다음처럼 정의될 수 있다.

max(+0.0, -0.0) = +0.0

min(+0.0, -0.0) = -0.0

하지만 단순 comparison-selection은 tie에서 어느 operand를 선택하는지에 따라 결과 부호가 달라질 수 있다.

예를 들어:

select(x >= z, x, z)

에서 x = -0.0, z = +0.0이면 equality가 true이므로 x를 선택할 수 있다.

result = -0.0

이는 maximum의 기대 의미인 +0.0과 다를 수 있다.

따라서 direct maximum/minimum instruction과 comparison-selection은 signed zero에서 다를 수 있다.

13. Infinity 처리

Infinity는 정상적인 순서 관계에 참여한다.

-Inf < finite < +Inf

따라서:

max(+Inf, x) = +Inf

min(-Inf, x) = -Inf

max(-Inf, x) = x

min(+Inf, x) = x

이 성질은 reduction 초기값으로 활용된다.

Max reduction 초기값:

acc = -Inf

Min reduction 초기값:

acc = +Inf

그 후:

acc = max(acc, x[i])

또는:

acc = min(acc, x[i])

로 갱신한다.

14. 정수 Maximum과 Minimum

정수에서는 NaN과 signed zero가 없기 때문에 의미가 비교적 단순하다.

max_int(x, z)

min_int(x, z)

그러나 signed와 unsigned 비교를 구분해야 한다.

동일한 bit pattern:

0xFFFFFFFF

Signed INT32:

-1

Unsigned UINT32:

4294967295

따라서:

signed max(-1, 1) = 1

unsigned max(4294967295, 1) = 4294967295

정수 max/min의 lowering에서는 signedness가 비교 조건에 반영되어야 한다.

15. Maximum과 Selection의 관계

Maximum은 다음처럼 표현할 수 있다.

p = x >= z
y = select(p, x, z)

Minimum은 다음과 같다.

p = x <= z
y = select(p, x, z)

추상 실행 구조:

Load X
→ Load Z
→ Compare
→ Predicate
→ Select
→ Store Y

하지만 direct max/min instruction이 있다면 다음처럼 축약된다.

Load X
→ Load Z
→ Max/Min
→ Store Y

이 변환은 다음 효과를 가질 수 있다.

Predicate register 제거 가능
Comparison instruction 제거
Selection instruction 제거
Instruction count 감소
Dependency chain 단축 가능

다만 특수값 semantics가 동일해야 한다.

16. Maximum과 Branch의 관계

Maximum을 branch로도 구현할 수 있다.

if x >= z:
    y = x
else:
    y = z

실행 구조:

Compare
→ Branch
→ Move X 또는 Move Z

Elementwise maximum처럼 경로가 매우 짧은 경우 branch는 일반적으로 불필요한 control-flow를 만들 수 있다.

GPU에서는 warp 내 thread별 조건이 다르면 divergence가 발생할 수 있다.

따라서 다음 방식이 더 자연스러울 수 있다.

Direct Max Instruction

또는:

Compare + Select

하지만 compiler가 실제로 어떤 형태를 선택하는지는 SASS를 확인해야 한다.

17. ReLU

ReLU는 scalar Maximum의 대표 사례다.

y[i] = max(x[i], 0)

수학적 의미:

x[i] > 0이면 y[i] = x[i]
x[i] <= 0이면 y[i] = 0

가능한 구현은 다음과 같다.

Direct Maximum

y = max(x, 0)

Comparison과 Selection

p = x > 0
y = select(p, x, 0)

Branch

if x > 0:
    y = x
else:
    y = 0

일반 유한값에서는 같은 결과를 만들 수 있지만 다음에서 차이가 있을 수 있다.

NaN
+0.0
-0.0
Fast-math
Direct maximum의 특수값 규칙

18. Clamp

Clamp는 Maximum과 Minimum의 합성이다.

t[i] = max(x[i], lower)
y[i] = min(t[i], upper)

하나의 식으로 쓰면:

y[i] = min(max(x[i], lower), upper)

실행 구조:

Load X
→ Max with Lower
→ Min with Upper
→ Store Y

또는:

Compare Lower
→ Select
→ Compare Upper
→ Select

lower <= upper이면 결과는 지정된 범위 안에 있다.

lower <= y[i] <= upper

Clamp는 다음 성질을 가진다.

clamp(clamp(x, lower, upper), lower, upper)
=
clamp(x, lower, upper)

즉 멱등적이다.

19. Saturation

Saturation은 출력값을 표현 가능한 범위에 제한하는 연산이다.

예를 들어 INT8 범위는 다음과 같다.

-128 <= y <= 127

Saturating conversion은 다음처럼 표현할 수 있다.

t = max(x, -128)
u = min(t, 127)
y = convert_to_int8(u)

일반 integer cast의 wraparound와 saturation은 다른 의미다.

wraparound:
범위를 넘으면 bit truncation 또는 modulo 의미

saturation:
범위를 넘으면 경계값으로 고정

Quantization에서는 이 차이가 중요하다.

20. Gradient Clipping

Elementwise gradient clipping은 Clamp로 표현할 수 있다.

clipped_grad[i]
=
min(max(grad[i], -limit), limit)

Maximum과 Minimum의 합성이다.

Norm-based gradient clipping은 전체 gradient norm reduction이 필요하므로 다른 구조다.

Elementwise clipping:
각 원소 독립

Norm clipping:
전체 norm reduction
→ 공통 scale
→ elementwise multiply

같은 “clipping”이라는 이름을 사용해도 실행 모티프가 다르다.

21. Max Reduction의 기반

Maximum 연산자가 결합성을 가지므로 여러 값을 tree 형태로 결합할 수 있다.

m = max(x0, x1, x2, x3)

순차 형태:

m0 = max(x0, x1)
m1 = max(m0, x2)
m2 = max(m1, x3)

Tree 형태:

m0 = max(x0, x1)
m1 = max(x2, x3)
m  = max(m0, m1)

GPU에서는 다음 구조로 확장된다.

thread-local max
→ warp max reduction
→ block max reduction
→ final max

따라서 Elementwise Maximum은 Max Reduction의 원자적 결합 연산자다.

22. Softmax에서의 Maximum

Stable Softmax는 row의 최댓값을 먼저 계산한다.

m = max_j(x[j])

그다음 모든 값에서 m을 뺀다.

e[j] = exp(x[j] - m)

Maximum은 수치 안정성을 위한 기준점을 제공한다.

x[j] - m <= 0

따라서:

exp(x[j] - m) <= 1

가 되어 exponential overflow 위험을 줄인다.

여기서 Maximum은 elementwise max가 아니라 reduction max다.

23. Argmax와의 관계

Argmax는 최대값뿐 아니라 해당 위치도 반환한다.

index =
argmax_i(x[i])

실행 상태는 다음 두 값을 함께 유지해야 한다.

current_max_value
current_max_index

갱신은 다음과 같다.

if x[i] > current_max:
    current_max = x[i]
    current_index = i

즉 Maximum과 Selection이 값과 index 두 곳에 동시에 적용된다.

Tie-breaking 규칙도 필요하다.

첫 번째 최대 index
마지막 최대 index
가장 작은 index
임의의 index

Maximum 자체보다 의미가 더 복잡하다.

24. 추상 실행 모티프

Tensor-Tensor Elementwise Maximum:

Index
→ Load X
→ Load Z
→ Maximum
→ Store Y

Scalar Maximum:

Index
→ Load X
→ Reuse Scalar
→ Maximum
→ Store Y

Compare-Select 구현:

Index
→ Load X
→ Load Z
→ Compare
→ Predicate
→ Select
→ Store Y

Clamp:

Index
→ Load X
→ Maximum with Lower
→ Minimum with Upper
→ Store Y

Reduction Maximum:

Load Multiple Values
→ Local Maximum Accumulation
→ Cross-thread Maximum Reduction
→ Final Store

25. 비용 구조

FP32 Tensor-Tensor Maximum의 원소당 기본 memory traffic은 다음과 같다.

x[i] load: 4 bytes
z[i] load: 4 bytes
y[i] store: 4 bytes

합계: 약 12 bytes

핵심 연산은 maximum 한 번이다.

1 comparison-selection 의미 / 12 bytes

산술 집약도가 낮으므로 독립 Elementwise Maximum kernel은 대체로 memory-bound가 되기 쉽다.

Scalar Maximum은 다음과 같다.

x[i] load: 4 bytes
y[i] store: 4 bytes

합계: 약 8 bytes

Scalar bound를 register나 immediate로 재사용할 수 있다.

26. 표현 계층별 형태

26.1 수학 계층

y[i] = max(x[i], z[i])

w[i] = min(x[i], z[i])

26.2 연산 그래프 계층

X ─┐
   ├→ Maximum → Y
Z ─┘

X ─┐
   ├→ Minimum → Y
Z ─┘

일부 그래프에서는 다음 이름을 사용할 수 있다.

Max
Min
Maximum
Minimum
Clip
Relu

26.3 Loop 계층

for (int i = 0; i < n; ++i) {
    y[i] = max(x[i], z[i]);
}

for (int i = 0; i < n; ++i) {
    y[i] = min(x[i], z[i]);
}

Comparison-Selection 형태:

y[i] = x[i] >= z[i] ? x[i] : z[i];

26.4 CUDA Kernel 계층

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

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

Minimum:

y[i] = fminf(x[i], z[i]);

비교 연산자와 ternary를 사용할 수도 있다.

y[i] =
    x[i] >= z[i] ? x[i] : z[i];

두 표현의 특수값 semantics는 반드시 비교해야 한다.

26.5 PTX 계층

개념적으로 다음과 같은 형태가 가능하다.

max.f32
min.f32

또는:

setp
selp

정수형에서는 signed/unsigned modifier가 포함될 수 있다.

26.6 SASS 계층

Architecture에 따라 다음 종류의 명령 또는 패턴을 예상할 수 있다.

FMAX
FMNMX
IMNMX
FSETP + FSEL
ISETP + SEL

정확한 mnemonic은 GPU architecture에 따라 달라질 수 있다.

핵심 데이터 흐름은 다음과 같다.

LDG X ─┐
       ├→ MAX/MIN → STG Y
LDG Z ─┘

27. SASS에서 Maximum과 Minimum을 식별하는 기준

27.1 두 입력 source

R_x ─┐
     ├→ MAX/MIN → R_y
R_z ─┘

27.2 단일 destination

더 크거나 작은 값 하나가 destination에 기록된다.

27.3 직접 명령 또는 Compare-Select

다음 두 패턴을 모두 확인해야 한다.

Direct:
MAX/MIN

Expanded:
Compare
→ Predicate
→ Select

27.4 Cross-thread communication 부재

Elementwise Maximum에는 일반적으로 다음이 없다.

Shuffle
Barrier
Shared memory reduction
Atomic operation

27.5 Reduction 문맥 구분

동일 register가 반복적으로 max accumulator로 사용되거나 shuffle과 함께 사용되면 max reduction일 가능성이 있다.

28. Direct Max/Min이 생성되지 않는 경우

특수값 semantics 차이

Compiler가 source expression의 NaN 규칙을 보존하기 위해 Compare-Select를 사용할 수 있다.

복합 predicate

단순 max/min이 아닌 추가 조건이 포함될 수 있다.

if valid:
    max(x, z)
else:
    default

데이터 타입

특정 dtype에서 native max/min 지원이 제한되거나 conversion이 필요할 수 있다.

관측 가능한 predicate

비교 결과가 다른 consumer에서도 사용되면 predicate를 별도로 유지할 수 있다.

Compiler option

Fast-math와 strict floating-point 설정에 따라 lowering이 달라질 수 있다.

29. Maximum과 Minimum의 Fusion

Maximum과 Minimum은 다른 elementwise 연산과 결합하기 쉽다.

예:

t[i] = a*x[i] + b
y[i] = max(t[i], 0)

분리 실행:

Load X
→ FMA
→ Store T

Load T
→ Max with 0
→ Store Y

융합 실행:

Load X
→ FMA
→ Max with 0
→ Store Y

중간 tensor의 store/load를 제거할 수 있다.

GEMM epilogue에서는 다음과 같다.

GEMM accumulator
→ Bias Add
→ Maximum with 0
→ Final Store

이는 GEMM + Bias + ReLU fusion이다.

30. In-place 실행

다음과 같이 입력 buffer를 갱신할 수 있다.

x[i] = max(x[i], lower)

x[i] = min(x[i], upper)

안전 조건은 다음과 같다.

기존 x가 다른 consumer에서 필요하지 않음
각 thread가 자신의 위치만 갱신함
Bound tensor와 output alias가 의미를 훼손하지 않음
Backward에서 원래 입력이 필요한지 확인됨

ReLU backward에서는 원래 입력 대신 출력의 양수 여부를 사용할 수 있는 구현도 있지만, 모든 activation에 동일하게 적용되지는 않는다.

31. 허용되는 변환

Operand 교환

max(x, z) → max(z, x)

min(x, z) → min(z, x)

특수값과 tie-breaking 의미를 확인해야 한다.

중복 제거

max(x, x) → x

min(x, x) → x

중첩 단순화

max(max(x, a), a)
→ max(x, a)

min(min(x, a), a)
→ min(x, a)

Compare-Select를 Direct Max/Min으로 변환

NaN 및 signed zero semantics가 일치할 때 가능하다.

Branchless 변환

짧은 branch를 direct max/min으로 바꿀 수 있다.

Clamp 결합

연속 lower/upper bound를 하나의 Clamp 모티프로 결합할 수 있다.

Producer/Consumer Fusion

앞선 affine 연산이나 후속 scale과 같은 연산과 결합할 수 있다.

Reduction Tree 변환

결합성을 이용해 max/min reduction의 tree 구조를 변경할 수 있다.

32. 제한되는 변환

NaN 규칙 무시

NaN 전파형과 숫자 우선형 maximum은 다른 연산이다.

Signed Zero 무시

max(+0.0, -0.0)

와:

min(+0.0, -0.0)

의 결과 부호가 중요할 수 있다.

Direct Max와 Compare-Select의 무조건적 동일시

Tie-breaking과 NaN 처리에서 달라질 수 있다.

Elementwise와 Reduction 혼동

같은 maximum primitive를 사용하지만 dependency 구조가 다르다.

Signed와 Unsigned Integer 혼동

같은 bit pattern의 순서가 달라질 수 있다.

Bound 순서 무시

Clamp에서:

lower > upper

이면 일반적인 범위 제한 의미가 깨지고 구현별 결과가 달라질 수 있다.

비관측성 확인 없이 중간값 제거

Maximum 결과가 여러 consumer에서 필요하면 materialization 제거가 불가능할 수 있다.

33. 정확성 검증

기본 reference는 명확한 특수값 규칙과 함께 정의해야 한다.

reference_max[i] =
    specified_max_semantics(x[i], z[i])

reference_min[i] =
    specified_min_semantics(x[i], z[i])

검증 항목:

일반 양수
일반 음수
동일값
+0.0
-0.0
NaN이 첫 번째 operand
NaN이 두 번째 operand
두 operand 모두 NaN
+Inf
-Inf
Subnormal
Signed integer
Unsigned integer
Broadcasting
In-place execution

수치값이 아니라 선택 결과가 중요하므로 다음을 함께 확인한다.

bitwise equality
mismatch_count
signed-zero bit
NaN 여부
NaN payload 필요 시 payload

34. 중요한 입력 패턴

일반 순서

x < z
x > z
x == z

Signed Zero

x = +0.0
z = -0.0

순서를 바꾼 경우도 테스트한다.

NaN

x = NaN
z = finite

x = finite
z = NaN

x = NaN
z = NaN

Infinity

x = +Inf
z = finite

x = -Inf
z = finite

Integer signedness

동일 bit pattern을 signed와 unsigned로 해석한다.

Bound 경계

x = lower
x = upper
x = lower - delta
x = upper + delta

35. 실험 설계

실험 1: FP32 Direct Maximum

y[i] = fmaxf(x[i], z[i])

목적:

PTX/SASS direct max instruction 확인
기본 load-max-store 구조 확인
NaN 처리 확인

실험 2: FP32 Direct Minimum

y[i] = fminf(x[i], z[i])

목적:

Minimum lowering 확인
Maximum과 instruction 구조 비교

실험 3: Compare-Select Maximum

y[i] =
    x[i] >= z[i] ? x[i] : z[i]

목적:

Direct maximum과 SASS 비교
NaN 및 signed zero 결과 비교
Predicate 생성 여부 확인

실험 4: Compare-Select Minimum

y[i] =
    x[i] <= z[i] ? x[i] : z[i]

목적:

Direct minimum과 의미 차이 확인

실험 5: ReLU Maximum

y[i] = max(x[i], 0)

목적:

Scalar zero 처리
Direct maximum, select, branch 비교 기반 마련
Signed zero와 NaN 확인

실험 6: Clamp

y[i] =
    min(max(x[i], lower), upper)

목적:

Max-Min instruction chain
Compare-Select chain과 비교
Branchless lowering 확인

실험 7: Max Reduction

m = max_i(x[i])

목적:

Elementwise maximum과 reduction 구조 비교
Warp shuffle/shared memory 확인
Reduction tree와 NaN 규칙 분석

실험 8: INT32 Maximum

y[i] = max(x[i], z[i])

목적:

Integer max instruction 확인
FP32 maximum과 비교

실험 9: UINT32 Maximum

목적:

Signed/unsigned 조건 차이
동일 bit pattern 결과 비교

실험 10: NaN Semantic Matrix

다음을 비교한다.

fmaxf(x, z)
x >= z ? x : z
x > z ? x : z

목적:

Operand 순서와 comparison relation의 영향 확인
Strict/fast-math 차이 확인

실험 11: Signed Zero Matrix

다음을 비교한다.

max(+0.0, -0.0)
max(-0.0, +0.0)
min(+0.0, -0.0)
min(-0.0, +0.0)

목적:

Tie-breaking과 operand order 확인
Bitwise 결과 기록

실험 12: FMA + ReLU Fusion

y[i] =
    max(a*x[i] + b, 0)

목적:

FFMA + FMAX 구조 확인
분리 kernel과 memory traffic 비교
Register dependency 분석

36. 분석 항목

의미 수준

Maximum 또는 Minimum
NaN 전파 규칙
Signed zero 규칙
Signed/unsigned
Broadcasting axis
Elementwise 또는 reduction

코드 수준

Direct library function
Comparison + ternary
Explicit branch
Clamp composition
Fusion 여부
In-place 여부

PTX 수준

Max/min 명령
Set-predicate
Select
Type modifier
Signedness
Rounding 및 NaN 관련 modifier

SASS 수준

FMAX 또는 max/min 계열 명령
Integer max/min 명령
FSETP/ISETP
FSEL/SEL
Global load/store
Predicate register
Reduction 반복 패턴

실행 수준

Kernel time
Effective bandwidth
Instruction count
Branch efficiency
Register count
Fusion 전후 memory traffic
Reduction latency

정확성 수준

Bitwise equality
NaN result
Signed zero bit
Operand order dependence
Mismatch count
Clamp boundary correctness

37. 연구에서의 의미

Maximum과 Minimum은 Comparison과 Selection이 하나의 직접 실행 primitive로 축약될 수 있음을 보여준다.

핵심 관계는 다음과 같다.

수학적 관계:
두 값의 순서 비교

논리적 표현:
Compare
→ Predicate
→ Select

ISA 표현:
Direct MAX/MIN instruction

이 연산은 다음 연구 질문을 명확하게 드러낸다.

같은 수학식이 같은 실행 의미를 보장하는가

일반 유한값에서는 direct maximum과 comparison-selection이 같아 보인다.

하지만 NaN과 signed zero에서는 다를 수 있다.

하나의 명령이 어떤 복합 의미를 포함하는가

Direct max/min 명령은 comparison과 selection을 동시에 포함한다.

결합성이 병렬 reduction을 어떻게 허용하는가

Maximum과 Minimum의 결합성은 tree reduction의 수학적 근거다.

원자적 primitive와 복합 연산의 관계

Maximum + Minimum
→ Clamp

Maximum with zero
→ ReLU

Repeated Maximum
→ Max Reduction

실행 모티프의 재사용

같은 maximum primitive가 elementwise activation, reduction, pooling, Softmax 안정화에 사용된다.

38. 최종 정의

Maximum은 다음과 같이 정의할 수 있다.

두 입력 중 정의된 순서 관계에 따라
더 큰 값을 출력하는 연산

Minimum은 다음과 같다.

두 입력 중 정의된 순서 관계에 따라
더 작은 값을 출력하는 연산

기본 수식:

y[i] = max(x[i], z[i])

w[i] = min(x[i], z[i])

논리적 분해:

Maximum:
Compare x >= z
→ Select x or z

Minimum:
Compare x <= z
→ Select x or z

기본 실행 구조:

Index
→ Load X
→ Load Z
→ Max/Min
→ Store Output

대표적인 최적화 가능성:

direct max/min instruction
comparison-selection fusion
branch 제거
scalar bound reuse
kernel fusion
in-place execution
Clamp composition
reduction tree 변환
중복 bound 제거

중요한 제한:

Floating-point Maximum과 Minimum은
NaN 처리, signed zero, tie-breaking 규칙에 따라
서로 다른 의미를 가질 수 있다.

따라서 direct max/min,
comparison-selection,
branch 구현을 동일하다고 보기 전에
특수값 semantics를 확인해야 한다.

'SASS_Probe' 카테고리의 다른 글

하드웨어 비종속적 Semantic Lowering 체계로의 확장 (0)	2026.06.24
기본 연산 의미 명세 08 - ReLU (1)	2026.06.23
기본 연산 의미 명세 06 - Selection (0)	2026.06.23
기본 연산 의미 명세 05 - Comparison (0)	2026.06.22
기본 연산 의미 명세 04 - Fused Multiply Add, FMA (0)	2026.06.22