기본 연산 의미 명세 01 - Elementwise Map

1. 연산 개요

Elementwise Map 은 입력을 구성하는 각 원소에 동일한 함수나 규칙을 독립적으로 적용하는 연산

• 입력 원소 하나
• 해당 원소에 대한 독립적인 변환
• 대응하는 출력 원소 하나

딥러닝 모델에서 가장 자주 등장하는 기본 실행 패턴 중 하나.

 

2. 수학적 정의

입력 텐서 X

입력값 하나를 출력 값 하나로 변환하는 함수 f 가 있을 때 Elementwise Map 은 다음과 같이 정의된다.

• Y = Map(f, X)
• y[i] = f(x[i])

모든 위치 i 에는 동일한 함수 f 가 적용된다.

각 출력은 같은 위치의 입력으로부터 독립적으로 계산된다.

 

3. 다중 입력 Elementwise Map

Elementwise Map 은 하나의 입력만을 사용할 필요는 없다.

출력 y[i] 는 일반적으로 x[j], y[j] 와 같이 다른 위치의 값에는 의존하지 않는다.

 

4. 입력과 출력의 의미

일반적인 단항 Elementwise Map 은 입력과 출력의 shape 이 같다.

하지만 데이터 타입까지 반드시 같을 필요는 없다. 

단순히 shape 보존이 아니라 각 출력 위치가 대응하는 입력 위치의 값만을 이용해 계산

 

5. 핵심 데이터 의존성

가장 중요한 성질은 원소별 독립성이다.

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

원소 사이를 연결하는 dependency 가 없다.

이 특성 때문에 모든 출력 원소를 동시에 계산할 수 있다. 

이것이 병렬 실행에 적합한 가장 직접적인 이유다.

 

6. Cross-thread communication 의 부재

순수한 Elementwise Map 에서는 다른 원소의 결과를 필요로 하지 않는다. 

수학적 의미만 놓고 보면 다음 구조가 필요하지 않다.

• warp shuffle
• shared memory 를 통한 값 교환
• block synchronization
• reduction
• atomic operation
• 다른 thread 의 계산 결과 대기

실제 생성된 코드에서 이러한 구조가 발견된다면 다음 가능성을 확인해야 한다. 

• 실제 연산이 순수한 Elementwise map 이 아닌가
• broadcasting 이나 lookup 이 포함되었는가
• layout transformation 이 함께 수행되는가
• 구현에 불필요한 synchronization 이 있는가
• 다른 연산과 fusion 된 상태인가

무엇이 나타나는가 뿐만 아니라 무엇이 나타날 필요가 없는가도 중요한 의미 정보

 

7. 의미 불변성

다른 구현으로 변경하더라도 반드시 유지해야 하는 조건은 다음과 같다

7.1 위치 대응 관계

출력 위치 i 는 대응하는 입력 위치 i 를 이용해 계산되어야 한다.

입력 위치를 임의로 변경하면 다른 연산이 된다.

 

7.2 함수 의미

각 원소에 적용되는 함수 f 의 의미가 보존되어야 한다.

 

7.3 원소 간 독립성

출력은 다른 위치의 입력이나 출력에 의존하지 않아야 한다. 

 

7.4 입력과 출력의 대응 개수

• 여러 입력을 하나로 합치는 reduction 이나
• 하나의 값을 여러 위치로 확장하는 broadcast 와 구분

 

8. 구조적 불변성

표현 계층이 달라져도 대체로 다음 실행 모티프를 유지한다.

• Index
• Load
• Local Transform
• Store

조금 더 구체적으로 표현하면 

• 처리할 출력 index 계산
• index 가 유효한 범위인지 확인
• 해당 위치의 입력 load
• thread-local 함수 계산
• 대응 위치에 출력 store

함수 f 가 복잡해지면 local transform 단계의 명령 수는 증가한다.

순수한 Elementwise Map 이라면 다음 구조는 원칙적으로 추가되지 않는다.

• many-to-one reduction
• reduction 결과 broadcast
• cross-thread communication
• block-wide synchronization

 

9. 함수 합성과 연산 융합

가장 중요한 대수적 성질은 함수 합성이다.

• y[i] = g(f(x[i]))

의 연산이 있을 때 두 번의 map 은 한 번의 map 으로 결합된다.

중간 텐서가 외부에서 사용되지 않는다면, 두 연산은 하나의 kernel 로 fusion 할 수 있다.

 

10. 중간값의 관측 가능성

연속 연산의 생각

분리된 kernel 에서의 실행 가능

중간값이 다음 조건에 해당한다면 materialization 이 필요할 수 있다.

 

11. Broadcasting 을 포함한 Elementwise 연산

shape 이 다른 값이 반복 사용되는 경우가 많다.

channel bias 의 경우 여러 batch 위치에서 반복 사용된다.

수학적으로는 위치별 덧셈이지만 실행 구조에서는 재사용이 존재한다.

Broadcst 가 포함되면 다음 최적화 가능성이 생긴다

• parameter register 유지
• constant memory 또는 cache 활용
• 반복 global load 감소
• vectorized parameter load

 

12. 연산 순서의 자유도

서로 다른 출력 원소는 독립적이므로 원소 처리 순서는 자유롭게 변경할 수 있다.

다음을 구분 필요

• 서로 다른 원소의 실행 순서 : 대체로 자유롭게 변경 가능
• 같은 원소에 적용되는 함수 순서 : 데이터 의존성과 대수적 성질에 의해 제한

 

13. 병렬 실행

Elementwise Map 은 출력 원소가 독립적이므로 GPU 에서는 thread 단위로 쉽게 병렬화할 수 잇다.

입력 크기가 thread 수보다 큰 경우 하나의 thread 가 여러 원소를 처리할 수 있다.

수학적 의미는 변하지 않지만 실행 스케줄이 변경된다.

 

14. 추상 비용 구조

수식은 단순히 다음과 같을 수 있다.

• y_i = f(x_i)

하지만 실제 실행 비용은 다음 세 부분으로 나뉜다.

• 전체 비용
  • index 계산 비용
  • memory 이동 비용
  • 함수 f 의 계산 비용

 

Index 계산 비용

• thread index 읽기
• block offset 계산
• global index 계산
• boundary 확인
• 주소 계산

 

Memory 비용

• 입력 global load
• 출력 global store

 

Function 비용

• FADD
• FMUL
• FFMA
• FMAX
• EXP
  reciprocal
• comparison
• selction

함수 f 가 매우 단순하면 arithmetic 보다 memory 와 index 계산의 비중이 더 클 수 있다.

 

15. 산술 집약도와 Memory-bound 특성

FP32 단항 Elementwise Map 을  생각

원소 하나를 처리할 때 일반적으로 다음 memory traffic 이 발생한다.

• 입력 load : 4 bytes
• 출력 store : 4 bytes
• 합계 : 최소 8 bytes

연산이 단순한 덧셈 하나라면

• 산술 연산 약 1회
• memory 이동 약 8 bytes

산술 집약도는 매우 낮다

• Arithmetic Intensity
• 산술 연산량 / memory 이동량

단순한 ReLU, Add, Scale 같은 연산은 대체로 memory bandwidth 와 kernel launch overhead 의 영향을 크게 받는다.

이것이 Elementwise fusion 이 중요한 이유

Fusion 은 산술 연산 자체를 크게 줄이지 않아도 global memory 왕복을 줄일 수 있다.

 

16. Kernel Fusion

세 개의 elementwise kernel 이 연속 실행된다고 하자

• z1_i = x_i + bias_i
• z2_i = relu(z1_i)
• y_i = scale_i * z2_i

분리 실행에서는 다음과 같은 memory 흐름이 발생한다.

• x load
  • bias add
  • z1 store
• z1 load
  • relu
  • z2 store
• z2 load
  • scale
  • y store

융합 시

• x load 
  • bias add
  • relu
  • scale
  • y store

줄어드는 것은

• z1 global store / load
• z2 global store / load
• kernel lauch 두 번
• 추가 index 와 address 계산

그러나 fusion 이 항상 유리한 것은 아님

과도한 fusion 은 다음 문제를 만들 수 있다.

• register pressure 증가
• instruction dependency chain 증가
• code size 증가
• occupancy 감소
• 다른 consumer 와 중간값을 공유하기 어려워짐

단순히 결합할 수 있는가가 아니라

memory traffic 감소가 register 와 실행 자원 증가보다 큰 이점을 만드는가?

 

17. In-place 실행

입력의 기존 값이 더 이상 필요하지 않다면 출력이 입력 buffer 를 덮어쓸 수 있다.

Elementwise Map 은 같은 위치의 값만 읽기 때문에 in-place 실행과 잘 맞는 연산이다.

하지만 graph 전체의 데이터 의존성을 확인하지 않고 적용하면 다른 consumer 가 필요로 하는 입력을 파괴할 수 있다. 

 

18. Vectorization

연속된 여러 원소를 하나의 vector load 와 store 로 처리할 수 있다.

Elementwise 연산의 의미는 원소별로 유지되지만 실제 명령에서는 여러 원소가 하나의 memory transaction 이나 vector register 단위로 처리될 수 있다.

 

19. 표현 계층별 형태

19.1 수학 계층

• y_i = f(x_i)

여기서는 입력과 출력의 의미 관계만 표현된다.

 

19.2 연산 그래프 계층

• x
  • Elementwise f
  • y

연속 연산은 다음과 같이 나타난다

• x
  • add
  • relu
  • scale
  • y

이 계층에서는 fusion 후보와 producer consumer 관계를 쉽게 관찰할 수 있다.

 

19.3 Loop 계층

이 표현에선느 반복 범위와 원소별 독립성이 드러난다. 

 

19.4 CUDA Kernel 계층

수학적 index i 는 GPU 의 block 과 thread index 를 이용해 계산된다

• i = blockIdx.x * blockDim.x + threadIdx.x

 

19.5 PTX 계층

일반적으로 다음 종류의 연산이 나타난다

• thread / block index 읽기
• inter multiply / add 를 통한 index 계산
• boundary comparison
• predicate 또는 branch
• global load
• 함수 f에 해당하는 산술 연산
• global store

 

19.6 SASS 계층

GPU architecture 와 compiler 에 따라 정확한 명령은 달라지지만, 기본 모티프는 다음과 같다

• special register 에서 thread 정보 읽기
  • index / address 계산
  • boundary predicate
  • global load
  • thread-local arithmetic
  • global store
  • exit

예상 가능한 명령 종류는 다음과 같다

• S2R
• IMAD
• IADD
• ISETP
• LDG
• FADD
• FMUL
• FFMA
• FMAX
• FMIN
• FSEL
• MUFU
• STG
• BRA
• EXIT

중요한 것은 특정 명령 이름보다 데이터 흐름이다

• input LDG
  • local register dependency chain
  • output STG
    
    

 

20. SASS 에서 Elementwise MAP 을 식별하는 기준

20.1 Load-Compute-Store 구조

가장 기본적인 signature 는 다음과 같다

• LDG input
  • local arithmetic
  • STG output

 

20.2 Cross-thread 명령 부재

순수한 Elementwise Map 에서는 일반적으로 다음 명령이나 구조가 없어야 한다.

• warp shuffl
• barrier
• shared memory reduction
• atomic accumulation

 

20.3 짧은 중간값 수명

입력에서 파생된 중간값은 한 thread 의 register 에서 생성되고 최종 store 이전까지만 유지된다

• input register
  • intermediate register
  • output register
  • STG

 

20.4 함수별 하위 signature

Elementwise Mapd 의 공통 구조 안에서 함수에 따라 산술 부분이 달라진다.

• Scale
  • FMUL
• Affine
  • FFMA
• ReLU
  • FMAX
• Clamp
  • FMAX, FMIN
• Sigmoid
  • negate
  • exponential
  • add
  • reciprocal

따라서 SASS 분석에서는 다음 두 수준을 구분할 수 있다.

상위 구조

• Elementwise Map

하위 함수 signature

• ReLU, Clamp, Affine, Sigmoid

 

21. 허용되는 변환

다음 변환 가능성이 존재한다

• 병렬화
• Vectorization
• Kernel fusion
• In-place execution
• Producer epilogue fusion
• Redundant operation 제거
• Parameter reuse

 

22. 제한되는 변환

• 함수 순서의 임의변경
• 위치 대응 관계 변경
• Broadcast domain 변경
• 관측 가능한 중간값 제거
• 부정확한 근사
• 비결정적 side effect 추가

 

23. 정확성 검증

Elementwise Map 은 reference CPU 구현과 직접 비교할 수 있다.

기본 검증 항목

• max_abs_diff
• max_relative_diff
• bitwise equality
• NaN 처리
• Inf 처리
• boundary input

함수에 따라 다음 입력을 포함해야 한다

• 0
• 양수
• 음수
• 매우 큰 값
• 매우 작은 값
• NaN
• +Inf
• -Inf
• threadhold 근처의 값

 

24. 실험 설계

Elementwise Map 의 기본 실험은 함수의 복잡도를 단계적으로 높이는 방식이 적절하다

실험 1 : Identity

y_i = x_i

순수 index/load/store 구조 확인

 

실험 2 : Scalar Add

y_i = x_i + c

FADD 추가 확인

 

실험 3 : Scalar Multiply

y_i = a * x_i

FMUL 확인

 

실험 4 : Affine

y_i = a * x_i + b

FMUL + FADD 가 FFMA 로 lowering되는 조건 확인

 

실험 5 : ReLU

y_i = max(0, x_i)

max 명령과 predicate / select 비교

 

실험 6 : Clamp

y_i = min(max(x_i, lower), upper)

두 경계 조건의 branchless lowering 확인

 

실험 7 : 연속 연산 분리와 Fusion

z_i = a * x_i + b

y_i = relu(z_i)

• 중간 global store / load 제거
• kernel 수 변화
• register dependency 변화

 

25. 연구에서의 의미

Elementwise Map 은 단순한 기초 연산이지만 연구 전체에서 중요한 기준점이 된다.

이 연산은 다음을 가장 명확하게 보여준다

• 수학적 독립성
• 함수 합성
• 비관측 중간값
• 낮은 산술 집약도

 

26. 최종 정의

핵심 의미는 다음 세 가지

• 위치별 대응
• 원소 간 독립성
• 동일한 local transform 의 반복

추상 실행 구조는 다음과 같다

• Index
• Load
• Local Transform
• Store

 

주요 최적화 가능성

• 병렬화
• vectorization
• kernel fusion
• in-place execution
• epilogue fusion
• 중간 materialization 제거