기본 연산 의미 명세 02

1. 연산 개요

Elementwise Add 는 두 입력의대응하는 원소를 더하여 출력하는 연산이다.

가장 기본적인 형태는 다음과 같다

y_i = x_i + z_i

각 출력, y_i 는 같은 위치에 있는 두 입력, x_i, z_i 에만 의존한다.

다른 위치의 입력 원소는 사용하지 않는다.

다음 두 성질을 동시에 가진다.

Elementwise Map 의 위치별 독립성
덧셈 연산의 대수적 성질

Elementwise Add 는 단순한 연산이지만 다음과 같은 복합 연산의 구성 요소로 반복적으로 등장한다.

Bias Add
Residual Add
Skip Connection
Gradient Accumulation
Tensor 합성
GEMM 또는 Convolution epilogue
LayerNorm 의 shift
Attention mask 적용
Optimizer 의 parameter update
여러 branch 의 출력 병합

2. 기본 수학적 정의

입 출력 shape 가 일반적으로 동일

3. Scaler Add

입력 텐서의 모든 원소에 동일한 상수 c 를 더하는 연산도 Elementwise Add 의 한 형태다

이 경우 상수 c 는 모든 출력 원소에서 재사용된다.

Scalar Add 에서는 두 번째 입력의 반복 global memory load 를 제거하거나 줄일 수 있다.

4. Broadcasting Add

실제 딥러닝에서는 두 입력의 shape 가 완전히 같지 않아도 broadcasting 을 통해 덧셈을 수행한다.

중요한 점은 broadcasting 이 단순한 덧셈의 의미에 추가적인 재사용 구조를 만든다는 것이다.

수학적 연산
- 대응되는 값의 덧셈
실행 구조
- 작은 입력의 반복적인 재사용

5. 입력과 출력 domain

입력과 출력이 같은 데이터 타입을 사용하는 경우가 많지만 반드시 그런 것은 아니다

6. 데이터 의존성

Elementwise Add 의 출력은 같은 위치의 두 입력에 의존한다.

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

따라서 모든 원소를 독립적으로 병렬 실행할 수 있다.

7. Elementwise Add 와 Reduction Add 의 구분

동일한 산술 연산이더라도 의미가 다르다.

Elementwise Add

y_i = x_i + z_i

구조

two values
one corresponding output

Sum Reduction

s = x_i + ...

구조

many values
one aggregate output

출력 위치별 독립성을 가진다.

Reduction 은 여러 입력 원소를 하나의 결과로 결합하므로 cross-element dependency 를 가진다.

따라서 같은 FADD 명령이 나타나더라도 전체 실행 모티프가 다르다.

개별 FADD 의 존재보다 데이터 흐름과 communication 구조를 함께 봐야한다.

8. 의미 불변성

Elementwise Add 를 다른 방식으로 구현하더라도 다음 조건은 보존되어야 한다

8.1 대응 위치의 보존

일반적인 Elementwise Add 에서는 같응 위치 관계가 유지되어야 한다.

다른 위치의 입력을 더하면 다른 연산이 된다.

8.2 두 입력에 대한 의존성

한 입력을 제거하면 일반적인 Elementwise Add 의 의미가 달라진다

8.3 출력 cardinality

입력의 각 논리적 위치마다 하나의 출력이 생성된다.

원소 수를 줄이거나 늘리지 않는다

8.4 Shape 및 broadcasting 규칙

입력 shape 이 다르다면 broadcasting 규칙이 정확하게 유지되어야 한다.

9. 대수적 성질

Elementwise Add 의 최적화 가능성은 덧셈의 대수적 성질과 밀접하게 연결된다.

9.1 교환 법칙

두 입력의 순서를 바꾸어도 수학적 결과는 같다

이 성질은 compiler 가 oerand 순서를 변경할 수 있는 근거가 된다.

9.2 결합 법칙

여러 elementwise Add 를 하나의 덧셈 tree 로 재배치할 수 있다.

부동소수점 반올림에 의해 bitwise equality 는 보장되지 않을 수 있다.

9.3 항등원

덧셈의 항등원은 0

이는 불필요한 Add 연산 제거의 수학적 근거가 된다 .

9.4 역원

덧셈의 역원은 -x

그러나 부동소수점에서는 문제 발생

+Inf + (-Inf) = NaN

이러한 문제들이 존재

9.5 선형성

Elementwise Add 는 선형 연산

또한 여러 선형 연산과 결합할 수 있다.

이 성질은 연산 재배치 가능성을 제공하지만, 실제로 어느 표현이 더 효율적인지는 메모리와 연산량에 따라 달라진다.

수학적 동등성이 곧 성능 최적화를 의미하는 것은 아니다

10. 정수 덧셈과 부동소수점 덧셈

Elementwise Add 는 데이터 타입에 따라 의미와 예외 조건이 달라진다

10.1 정수 덧셈

고정 비트폭 정수에서는 overflow 가 발생할 수 있다.

10.2 부동소수점 덧셈

부동소수점 덧셈은 개념적으로 다음 과정을 거친다

두 값의 exponent 정렬
significand 덧셈
normalization
목표 precision 에 맞게 반올림

따라서 실수 덧셈과 달리 다음 현상이 발생

작은 값이 큰 값에 더해져도 결과가 변하지 않을 수 있음

small 이 수학적으로 0 이기 때문이 아니라, 제한된 precision 에서 반올림되기 때문

11. 부동소수점에서의 비결합성

세 값이 존재

a : 매우 큰 양수
b : 매우 큰 음수
c : 작은 양수

연산 순서에 따라 결과가 c 와 0 으로 달라질 수 있다.

이는 여러 Add 를 fusion 하거나 residual branch 를 재배치할 때 중요해진다.

연산 순서 변경 시 다음 구분

실수 수학에서의 동등성
부등소수점 수치 오차
bitwise 재현성
모델 수준의 의미 보존

12. 특수값 처리

일반적인 딥러닝에서는 signed zero 차이가 중요하지 않은 경우가 많지만, bitwise equivalence를 평가한다면 고려해야 한다.

13. 추상 실행 모티프

Tensor-Tensor elementwise add 의 기본 실행 모티프

현재 thread 가 처리할 index 계산
boundary 확인
x_i 주소 계산
z_i 주소 계산
x_i load
z_i load
두 값을 더함
y_i 주소 계산
결과 store

14. 비용 구조

전체 비용
- index 및 address 계산
- 입력 memory load
- 덧셈 연산
- 출력 memory store

FP32 Tensor-Tensor Add 에서 원소 하나당 최소 memory traffic 은 대략 다음과 같다

x_i load : 4 bytes
z_i load : 4 bytes
y_i store : 4 bytes

산술 연산은 일반적으로 덧셈 1 회

1 floting-point operation / 12 bytes

따라서 산술 집약도가 매우 낮다

Elementwise Add 는 일반적으로 compute-bound 보다 memory-bound 되기 쉽다

15. Scalar Add 의 비용 특성

Scaler Add 에서는 상수 c 를 매 원소마다 global memory 에서 읽을 필요가 없다.

상수는 다음 방식 중 하나로 전달될 수 있다.

kernel argument
register
constant memory
immediate or encoded constant
uniform register

16. Broadcasting Add 의 비용 특성

작은 입력이 여러 출력에서 반복 사용된다.

논리적으로는 출력 원소마다 bias 를 하나씩 읽지만, 실제 hardware cache 나 register reuse 에 의해 memory traffic 을 줄일 수 있다.

17. 표현 계층별 형태

17.1 수학 계층

대응 원소의 합이라는 의미만 나타난다.

17.2 연산 그래프 계층

두 입력의 producer
출력의 consumer
broadcasting 관계
fusion 후보
중간 텐서의관측 가능성

17.3 Loop계층

동일 shape
Scalar Add
Braodcasting Add

17.4 CUDA Kernle 계층

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

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


>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>

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

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

17.5 PTX 계층

다음 종류의 연산 예상 가능

thread / block index read
global index cal
boundary predicate
input X global load
input Z global load
floating-point add
output global store

핵심 산술 명령은 개념적으로 다음과 같다

add.f32

하지만 compiler 최적화나 주변 연사에 따라 Add 가 다른 명령에 흡수될 수 있다.

17.6 SASS 계층

독립된 FP32 Add kernel 에서는 일반적으로 다음 모티프를 예상할 수 있다.

S2R / index 관련 명령
IMAD or IADD
ISETP
LDG X
LDG Z
FADD
STG Y
EXIT

18. Add 가 독립된 FADD 로 나타나지 않는 경우

18.1 FMA 로 결합

compiler 는 곱셈과 덧셈을 결합할 수 있다.

18.2 상수 folding

덧셈 항등원이 더해질 경우 Add 명령의 제거 가능

18.3 연속 상수 결합

+1+2 의 경우

+3 과 같이 단순화 가능

18.4 Producer or consumer and fusion

GEMM 결과에 bias 를 더하는 경우

분리된 kernel 이라면 FADD 가 나타날 수 있지만, epilogue fusion 이 적용되면 accumulator 의 최종 store 전에 Add 가 수행된다.

이때 SASS 에서이ㅡ위치와 register lifetime 이 크게 달라진다.

19. SASS 에서 Elementwise Add 를 식별하는 기준

19.1 두 입력 load

Tensor-Tensor Add 에서는 일반적으로 두 개의 입력 경로가 존재

19.2 하나의 local Add

두 입력 register 가 하나의덧셈 명령으로 결합된다.

FAD R_out, R_x, R_z

19.3 하나의 출력 store

STG ooutput

19.4 Cross-thread communication 부재

순수한 Elementwise Add 에는 일반적으로 다음이 필요하지 않다.

SHFL
BAR
shared memory reduction
atomic add
다른 thread 결과 대기

19.5 짧은 dependecy chain

LDG
FADD
STG

산술 dependency chain 이 짧다.

연산 자체보다 memory latency 와 throughput 이 성능에 더 큰 영향을 줄 수 있다.

20. Atomic Add 와의 구분

Atomic Add 는 단순 Elementwise Add 의 원소별 독립성을 가지지 않는다.

21. In-place Add

입력 하나에 결과를 덮어쓸 수 있음

별도 출력 buffer 가 필요하지 않으므로 memory 사용량을 줄일 수 ㅣㅇㅆ지만 global memory traffic 자체는 대체로 다음과 같이 남는다.

x load
z load
x store

22. Residual Add

Elementwise Add 의 대표적인 고수준 사례

수학적으로는 일반적인 Tensor-Tensor Add 와 같다

그러나 모델 구조에서는 두 입력의 의미가 다르다.

main : 반환된 주 경로
residual : 이전 표현을 우회해서 전달한 경로

Residual Add 의 출력이 곧바로 nomalization 에 사용된다면 중간 materialization 제거 검토 가능

하지만 nomalization 은 reduction 을 포함하므로 단순 Elementwise Add fusion 보다 복잡

23. Bias Add

broadcasting 을 포함한 Elementwise Add

producer epilogue 로 fusion 되는 대표 사례

24. Maks Add

단순 덧셈이지만, mask 의 의미와 broadcasting domain 이 반드시 보존되어야 한다.

25. Grandient Accumulation

수학적으로는 Elementwise Add 와 동일하지만

gradient 가 여러 kernel 에서 비동기적으로 누적된다면 Atomic Add 가 필요할 수 있다.

26. Kernel Fusion

Elementwise Add 는 앞뒤 연산과 fusion 하기 쉽다.

27. 연속 Add 의 결합

부동 소수점에서는 결과가 약간 달라질 수 있음

28. Vectorization

연속된 여러 원소를 vector 단위로 load / store 가능

29. 연구에서의 의미

단순히 GPU 가 덧셈 명령 하나를 실행하는 사례가 아니다.

원소별 독립성
- thread-level 병렬화
두 입력의 대응 관계
- 두 개의 load 경로
낮은 산술 집약도
- memory-bound 실행
항등원과 상수 결합
- 불필요한 명령 제거
함수 합성과 비관측 중간값
- kernel fusion
Multiplt + Add
- FFMA lowering
실수 덧셈의 결합 법칙
- 실행 tree 재배치 가능성
부동소수점 비결합성
- bitwise 결과와 수치 동등성의 구분

'SASS_Probe' 카테고리의 다른 글

기본 연산 의미 명세 04 - Fused Multiply Add, FMA (0)	2026.06.22
기본 연산 의미 명세 03 - Elementwise Multiply (0)	2026.06.22
기본 연산 의미 명세 01 - Elementwise Map (0)	2026.06.21
연산 의미 명세와 불변성 기반 실행 구조 분석 방법론 (0)	2026.06.21
수학적 연산의 SASS Lowering 과 불변성 기반 GPU 최적화 탐구 (0)	2026.06.20

기본 연산 의미 명세 02 - Elementwise Add