Execution Primitive Lap - operator 에 반복적으로 나타나는 구조적 실행 특징들을 ptimitive 단위로 분해하고, 이를 GPU 특성에 맞는 realization 으로 미리 구현, 검증하는 공간이다.

1. 정의

여기서의 primitive 는 단순한 scaler operation 이 아니라, 여러 operator 구현에서 반복적으로 재사용될 수 있는 실행 구조 단위를 뜻한다.

즉, 이 공간은 특정 operator 하나를 직접 구현하는 곳이 아니라, 이후 다양ㅇ한 operator 를 GPU 친화적으로 실현하기 위해 필요한 중간 실행 구성 요소를 분리하고 실험하는 곳이다.

 

2. 왜 필요한가

GPU 위에서의 kernel generation 은 단순히 수식을 코드로 옮기는 작업이 아니다. 

같은 수학적 의미를 가지는 계산이라도, 실제 GPU 실행에서는 다음 요소들에 따라 전혀 다른 성능과 비용 구조를 가진다.

• memory hierarchy 활용 방식
• tile / block / warp 단위의 병렬 조직
• synchronization 비용
• register / shared memory pressure
• data reuse 가능성
• recompute 와 materilization 사이의 trade-off
• floating - point order 변화에 따른 numerical behavior

임의의 operator 를 곧바로 하드웨어 특화 커널로 생성하려면, 먼저 그 operator 내부에 반복적으로 나타나는 구조를 더 작은 실행 단위로 분해할 수 있어야 한다.

 

즉, 이 프로젝트의 목적은 다음과 같다.

• operator 를 실행 관점의 구조적 특징으로 나눈다.
• 각 특징을 GPU 친화적인 primitive 로 대응시킨다
• primitive 별 realization 을 실제로 구현한다.
• GPU 조건에 따라 어떤 realization 이 유리한지 검증한다.
• 이후 kernel generation 에서 재사용 가능한 basis 를 만든다.

 

3. 이 프로젝트가 다루는 대상

이 프로젝트는 작은 연산을 다루는 공간이 아니다.

중심은 계산을 어떻게 조직하는가에 있다.

예를 들어 아래와 같은 것들이 대상이 된다.

reduction 계열

• tree reduction
• warp-local reduction
• block reduction
• hierarchical accumulation

streaming 계열

• online accumulation
• streaming weighted reduction
• blockwise state update
• rescaling-preserving accumulation

tile staging 계열

• shared-memory staging
• double buffering
• producer-consumer tile pipeline
• cp.async 기반 stage loading

rematerialization 계열

• intermediate save 제거
• recompute-based epilogue
• checkpoint-like realization
• memory-for-compute trade-off 구조

fusion / boundary hading 계열

• elementwise boundary elimination
• parial epilogue fusion
• accumulation-local fusion
• materialization 회피 구조

즉 이 프로젝트는 operator 보다 아래에 있으면서도, scalar op 보다는 위에 있는 실행 구조 계층을 다룬다.

 

4. 여기서 말하는 구조적 실행 특징이란 무엇인가

operator 마다 표현은 다르지만, 실제 구현 단계에서 반복적으로 등장하는 공통 구조가 있다.

이 문맥에서 구조적 실행 특징은 그런 공통 구조를 가리킨다.

예를 들면 다음과 같다

• 결과가 reduction 형태로 요약 가능한지
• partial state 를 유지하며 streaming update 가 가능한지
• tile 단위로 locality 를 만들 수 있는지
• intermediate 를 저장하지 않고 재구성할 수 있는지
• 순서 재배치가 성능상 유리한지
• blockwise decomposition 이 가능한지
• mergeable state 를 정의할 수 있는지

이런 특징들은 특정 operator 하나에 묶이지 않는다.

예를 들어  softamx, layernorm, attention 일부 statistics 연산은 서로 다르지만, online accumulation 이나 mergeable state 와 같은 공통 primitive 관점에서 다시 묶일 수 있다.

핵심은 개별 operator 이름이 아니라, 여러 operator realization 에 반복 등장하는 실행 구조를 분리하는 것이다.

 

5. 프로젝트의 역할

크게 네 가지

5.1 구조적 실행 특징을 primitive 로 분해

operator 를 직접 다루지 않고, 반복 가능한 실행 구조로 분해한다.

예

• softmax 전체를 바로 구현한느 대신, max-tracking, rescaled accumulation, weighted reduction 같ㅇ느 primitive 관점으로 나눔
• normalization 전체를 바로 구현하는 대신 online statistics, mergeable summary state, normalize stage 등으로 나눔

 

5.2 primitive 별 realization 구현

같은 primitive 라도 구현 방식은 여러 가지일 수 있다.

예

• reduction
  • warp shuffle 기반 / shared memory 기반 / multi-stage 기반
• tile staging
  • naive shared load / double buffer / cp.async stage pipeline
• rematerialization
  • saved intermediate 기반 / recompute 기반

즉, 이 프로젝트는 primitive 당 하나의 정답 구현을 만드는 공간이 아니라 realization family 를 구현하고 비교하는 공간이다.

 

5.3 hardware-aware 검증

primitive realization 은 GPU 구조와 분리될 수 없다.

같은 primitive 라도 다음 조건에 따라 유리한 구현이 달라진다.

• shared memory 크기
• register pressure
• occupancy
• memory transaction behavior
• warp scheduling
• architecture - specific async copy 지원 여부
• cache / shared reuse 특성

따라서 이 프로젝트는 단순 functional correctness 만 보는 것이 아닌, 

각 realization 이 실제 하드웨어에서 어떤 비용 구조를 가지는지 검증해야 한다.

 

5.4 code generation 을 위한 기반 제공

최종 목적은 ptimitive 구현 자체가 아니라

이 primitive 들을 future kernel generation 의 building block 으로 사용하는 것이다.

즉 나중에는 다음 흐름으로 이어진다.

Operator
→ structural execution features
→ execution primitives
→ GPU-specific realizations
→ hardware-aware kernel generation

이 프로젝트는 이 흐름에서

operator 와 final generated kernel 사이의 재사용 가능한 중간 계층을 만든다.

 

6. 다른 프로젝트와의 경계

6.1 GPU Probing Lab 과의 차이

gpu-probing-lab 은 GPU 의 실제 하드웨어 반응을 측정하고 역추적하는 공간이다.

• cache / bandwidth / latency 반응
• shared memory bank conflict
• tile shape 에 따른 hardware response
• staging depth 에 따른  latency hiding
• stride / coalischig behavior

반면 execution - primitive - lab 은

그 하드웨어 위에서 어떤 실행 구조를 사용할 수 있는지를 구현하고 비교한다.

즉

• gpu-probing-lab
  • gpu 가 어떻게 반응하는가
• execution-primitive-lab
  • 그 반응을 구려해 어떤 ptimitive realization 이 유효한가.

6.2 Property Atlas 와의 차이

property atlas 는 연산이 갖는 구조적 성질 자체를 정리하는 공간이다.

예

• associativity
• shift invariance
• mergeable statistics
• streaming 가능성
• rematerializable intermediate

반면 execution-primitive-lab 은

그 성질을 바탕으로 실제 실행 구조를 만든다.

즉,

• property atlas
  • 무엇이 가능한가
• execution-primitive-lab
  • 그 가능성을 어떻게 GPU 실행 구조로 실현하는가

 

6.3 Preservation Guide 와의 차이

preservation guide 는 어떤 성질을 보존하면서 어떤 변환이 허용되는지를 다루는 공간이다.

예

• 어떤 reorder 가 의미 보존적인가
• 어떤 state 만 유지하면 streaming 변환이 가능한가
• 무엇을 저장하지 않고 재구성해도 되는가

반면 execution-primitive-lab 은

그 보존 규칙을 바탕으로 실제 primitive realization 을 구현한다

즉,

• preservation guide
  • 무엇을 깨지 않고 바꿀 수 있는가
• execution-primitive-lab
  • 그 변환을 실제로 어떻게 구현할 것인가

 

7. 수학적 원리와 GPU 구현의 관계

수학적 원리 자체를 증명하는 공간이 아님

근본적 수학 원리는 GPU 와 독립적으로 존재

예

• associativity
• commutativity
• rescaling equivalence
• mergeable sufficient state
• decomposition 가능성

하지만 이 원리들은 GPU 위에 올라가는 순간 다음 질문으로 바뀐다.

• 이 원리들은 실제 병렬 구조로 바뀔 수 있는가
• 그 과정에서 어떤 memory / ssynchronization cost 가 발생하는가
• floating-point order 변화로 수치적 차이가 생기는가
• 어떤 hardware configuration 에서 실제로 유리한가

즉 execution-primitive-lab 은 수학적 원리 자체를 다루기보다, 그 원리가 GPU 위에서 어떤 realization 으로 바뀌는지 관찰하고 구현하는 공간이다.

 

8. 예시 : softamx 와 normalization 을 어떻게 볼 것인가

operator 이름 그대로 다루는 것이 아닌

중심은 그 안에 있는 구조적 특징이다.

softmax에서 볼 수 있는 것

• max tracking
• rescaled accumulation
• denominator coupling
• weighted reduction
• blockwise / streaming update 가능성

normalization 에서 볼 수 있는 것

• statistics accumulation
• mergeable summary state
• blockwise reduction
• online update 가능성
• normalize stage 분리 가능성

완성 operator 라기 보다 여러 primitive 가 조합되어 나타나는 사례로 다뤄진다.

 

9. 결과적으로 이 프로젝트가 만드는 것

단순한 커널 모음집이 아닌

9.1 Primitive catalog

반복적으로 등장하는 실행 primitive 의 분류와 정의

9.2 Realization family

primitive 별로 가능한 구현 방식들의 집합

9.3 Cost evidence

각 realization 이 어떤 GPU 조건에서 어떤 비용 구조를 보이는지에 대한 실험 결과

9.4 Codegen basis

향후 hardware-aware kernel generation 에 사용할 수 있는 재사용 가능한 구성 요소 집합

 

즉 산출물이 커널 하나가 아니라, 커널 생성을 위한 구조적 실행 라이브러리에 가깝다.

 

10. 핵심 문장 요약

operator 를 직접 구현하는 공간이 아니고, 단순한 low-level op 저장소도 아니다.

execution primitive lab 은 다음을 목표로 한다.

• operator 에 반복적으로 나타나는 구조적 실행 특징들을 분리하고
• 이를 primitive 단위로 재정의하며
• 각 primitive 를 GPU 특성에 맞게 realization 으로 구현, 검증하고
• 이후 hardware--aware kernel generation 에 재사용 가능한 basis 를 제공한다.