GPU Probing 결과를 코드 생성 규칙으로 연결하는 구조

1. 문제의식

GPU probing을 진행하다 보면 실험 범위는 자연스럽게 계속 확장된다.

처음에는 warp 단위의 실행 차이를 관찰하더라도, 곧 block 내부 warp 간 상호작용, shared memory, occupancy, SM scheduling, L2 cache, memory controller, grid-level scheduling까지 관심이 넓어진다.

하지만 이 프로젝트의 목적은 GPU microarchitecture 자체를 완전히 설명하는 것이 아니다.

이 프로젝트의 본래 목적은 다음에 가깝다.

AI compiler가 특정 연산 또는 layer에 대해 더 적절한 CUDA kernel code를 생성하기 위한 근거를 얻는 것.

따라서 GPU probing은 독립적인 하드웨어 탐사 프로젝트가 아니라, compiler의 code generation decision을 뒷받침하는 empirical evidence 수집 과정이어야 한다.

---

2. 핵심 관점

GPU probing의 결과는 단순히 다음과 같은 형태로 쓰이면 안 된다.

add가 mul보다 몇 % 빠르다.
shared memory load가 global memory load보다 몇 배 빠르다.
warp A가 warp B보다 progress가 높다.

이런 수치는 흥미롭지만, 바로 코드 생성 규칙으로 연결되기 어렵다.

더 중요한 것은 다음이다.

어떤 연산 구조가
어떤 실행 signature를 만들고,
그 signature가 어떤 병목 자원과 연결되며,
그 병목을 줄이기 위해 어떤 kernel variant를 선택해야 하는가.

즉 필요한 구조는 다음과 같다.

Probe Result
→ Execution Signature
→ Bottleneck Classification
→ Cost Model Signal
→ Codegen Decision Rule

---

3. 단순 연산 속도 비교의 한계

예를 들어, 더하기 연산이 곱하기 연산보다 몇 퍼센트 빠르다는 수치를 얻었다고 하자.

직관적으로는 이 비율에 맞춰 fused kernel 내부의 연산 배치를 조절하고 싶을 수 있다. 일종의 최소공배수처럼 서로 다른 latency를 가진 연산들을 비율에 맞춰 섞어 pipeline이 놀지 않게 만드는 발상이다.

하지만 GPU kernel code generation에서 실제 병목은 단순히 add와 mul의 순수 latency 차이로 결정되지 않는 경우가 많다.

실제 fused kernel에서 더 중요한 요소는 다음과 같다.

global memory load/store 비용
memory coalescing 여부
register dependency chain
instruction-level parallelism
warp-level latency hiding
shared memory bank conflict
synchronization cost
register pressure
occupancy 감소
warp divergence

따라서 add vs mul 같은 개별 연산 속도표는 cost model의 일부 신호가 될 수는 있지만, codegen rule의 중심 기준이 되기는 어렵다.

---

4. “최소공배수 느낌”의 정확한 해석

초기 직관은 버릴 필요가 없다. 다만 표현을 바꾸는 것이 좋다.

원래 직관:

서로 다른 연산 latency 비율에 맞춰 코드를 생성한다.

더 정확한 해석:

서로 다른 실행 자원의 latency와 throughput을 고려해
stall이 최소화되도록 독립 작업을 interleave하고,
전체 실행 구조의 balance를 맞춘다.

이것은 다음 개념에 가깝다.

latency-balanced scheduling
throughput-aware code generation
latency hiding
instruction interleaving
probe-driven cost model
kernel variant selection

즉, 핵심은 “연산 비율을 맞추는 것”이 아니라 “병목 자원이 놀지 않도록 실행 구조를 배치하는 것”이다.

---

5. Probing 결과가 가져야 하는 형태

좋지 않은 형태의 결과:

add는 mul보다 8% 빠르다.
global load는 shared load보다 느리다.

이런 결과는 하드웨어 관찰로는 의미가 있지만, compiler decision으로 직접 연결되기 어렵다.

좋은 형태의 결과:

elementwise fusion depth 1 → 100 GB/s
fusion depth 2 → 160 GB/s
fusion depth 4 → 210 GB/s
fusion depth 8 → 205 GB/s
fusion depth 16 → 170 GB/s

이 경우 compiler rule이 바로 도출된다.

elementwise chain은 4~8개 정도까지 fusion하는 것이 유리하다.
그 이상은 register pressure 또는 instruction pressure로 이득이 줄어든다.

즉, probing 결과는 단순한 성능 수치가 아니라 다음 정보를 포함해야 한다.

측정한 kernel variant
변화시킨 구조적 변수
성능 변화
병목 추정
codegen rule로의 환원

---

6. Probe Result Schema 제안

각 실험 결과는 다음과 같은 구조로 정리하는 것이 좋다.

const probeResult = {
  id: "elementwise_fusion_depth_probe",

  purpose: "elementwise op chain의 fusion depth가 throughput, register pressure, occupancy에 주는 영향을 측정한다.",

  measuredVariables: {
    fusionDepth: [1, 2, 4, 8, 16],
    bandwidthGBps: [100, 160, 210, 205, 170],
    registersPerThread: [16, 20, 28, 40, 72],
    occupancy: [1.0, 1.0, 0.75, 0.5, 0.25]
  },

  observation: {
    saturationPoint: 4,
    degradationPoint: 16,
    likelyBottleneck: "register_pressure"
  },

  executionSignature: {
    type: "memory_bound_to_register_pressure_transition",
    summary: "fusion depth가 증가할수록 global memory traffic은 줄지만, 일정 지점 이후 register pressure와 occupancy 감소로 throughput이 하락한다."
  },

  codegenRule: {
    targetPattern: "elementwise_chain",
    rule: "fuse_until_register_pressure_exceeds_threshold",
    preferredFusionDepth: "4_to_8_ops",
    avoid: "very_deep_fusion_without_reuse"
  }
};

이 구조의 핵심은 실험 결과가 codegenRule까지 도달한다는 점이다.

---

7. Compiler Cost Model의 기본 형태

probe 결과는 compiler 내부에서 일종의 empirical cost model로 쓰일 수 있다.

단순 모델은 다음과 같이 생각할 수 있다.

cost(kernel_variant)
=
memory_cost
+ compute_cost
+ dependency_cost
+ sync_cost
+ occupancy_penalty
+ register_pressure_penalty
+ bank_conflict_penalty

중요한 점은 cost model이 절대적으로 정확할 필요는 없다는 것이다.

초기 단계에서는 다음 정도만 가능해도 충분하다.

Variant A보다 Variant B가 유리한 조건을 구분한다.
특정 구조가 병목으로 바뀌는 지점을 찾는다.
fusion / tiling / reduction / shared memory 사용 여부를 결정한다.

즉, 목표는 완전한 성능 예측기가 아니라, code generation decision을 제한하고 유도하는 경험적 판단 모델이다.

---

8. Code Generation 흐름

최종적으로 compiler는 다음 흐름을 갖는 것이 이상적이다.

Operation Pattern 인식
↓
Candidate Kernel Variants 생성
↓
Probe 기반 Cost Model로 평가
↓
가장 낮은 예상 cost의 variant 선택
↓
CUDA Kernel Template Instantiation

예를 들어 LayerNorm은 다음과 같이 처리될 수 있다.

LayerNorm
↓
row-wise reduction pattern 인식
↓
Variant A: one warp per row
Variant B: one block per row
Variant C: two-pass reduction
↓
row length, dtype, alignment, shared memory usage 기준으로 cost 평가
↓
적절한 variant 선택
↓
blockDim, vectorized load, warp shuffle, shared memory partial reduction 설정

이때 GPU probing의 역할은 Variant A/B/C 중 무엇을 고를 것인가에 대한 근거를 제공하는 것이다.

---

9. 주요 실험 범위와 Codegen 연결

9.1 Elementwise Fusion Probe

관찰 대상:

fusion depth
global memory traffic
register count
occupancy
throughput

codegen rule:

중간 tensor write/read 비용이 큰 elementwise chain은 fusion한다.
단, register pressure가 threshold를 넘으면 fusion을 끊는다.

예상 규칙:

fusion depth 4~8까지는 적극적으로 fusion
fusion depth 16 이상은 cost model 확인 없이는 제한
reuse 없는 deep fusion은 피함

---

9.2 Memory Access Probe

관찰 대상:

contiguous access
stride access
coalesced / uncoalesced access
vectorized load
alignment
reuse distance

codegen rule:

contiguous axis를 threadIdx.x에 매핑한다.
stride가 큰 axis는 inner mapping에서 피한다.
alignment가 맞으면 vectorized load/store를 사용한다.

예상 규칙:

float4 / half2 / int4 가능한 경우 vectorized memory path 우선 고려
non-coalesced access가 강하면 layout transform 또는 axis reorder 검토

---

9.3 Warp Execution Signature Probe

관찰 대상:

independent ALU
dependent ALU chain
shared memory load
global memory load
mixed workload

codegen rule:

dependency chain이 긴 연산은 per-thread multiple elements 또는 instruction interleaving으로 ILP를 늘린다.
memory latency가 긴 구조는 충분한 ready warp supply를 확보한다.

예상 규칙:

dependent chain 중심 kernel은 unroll만 늘린다고 좋아지지 않는다.
독립 작업을 interleave할 수 있는 형태로 codegen한다.

---

9.4 Latency Hiding Probe

관찰 대상:

active warps per block
ready warp supply
memory-heavy warp progress
compute warp supply
stall reduction ratio

codegen rule:

memory-heavy kernel은 block당 warp 수를 너무 작게 잡지 않는다.
load 이후 바로 dependency chain으로 이어지는 코드를 피한다.
여러 element의 load와 compute를 interleave한다.

예상 규칙:

memory-bound kernel에서는 occupancy와 ILP 확보가 우선이다.
compute-heavy kernel에서는 register pressure와 instruction throughput을 더 우선한다.

---

9.5 Reduction Probe

관찰 대상:

row length
warp shuffle reduction
shared memory reduction
block-level reduction
sync cost
multi-pass reduction

codegen rule:

row length가 작으면 one-warp-per-row를 사용한다.
row length가 중간 크기면 one-block-per-row와 shared partial reduction을 사용한다.
row length가 매우 크면 multi-block 또는 two-pass reduction을 고려한다.

예상 규칙:

row length <= 32 → one warp per row
32 < row length <= 1024 → one block per row
very large row → multi-pass reduction

이 규칙은 Softmax, LayerNorm, ReduceSum에 직접 연결된다.

---

9.6 Shared Memory Trade-off Probe

관찰 대상:

global memory traffic reduction
shared memory latency
bank conflict
sync overhead
occupancy loss

codegen rule:

reuse가 충분하면 shared memory staging을 사용한다.
reuse가 낮으면 direct global load/store를 유지한다.
bank conflict가 크면 padding 또는 layout 변경을 적용한다.

예상 규칙:

GEMM tile staging → shared memory 사용
단순 one-pass elementwise → shared memory 사용하지 않음
row-wise reduction → row size와 reuse에 따라 선택
bank conflict 발생 layout → padding 삽입

---

10. Codegen Rule로 환원되는 예시

예시 1. Elementwise Chain

관찰:

fusion depth가 증가할수록 global memory traffic 감소
하지만 일정 depth 이후 register pressure 증가로 occupancy 감소

해석:

초기에는 memory-bound 개선 효과가 크다.
후반에는 register pressure가 새로운 병목이 된다.

codegen rule:

elementwise chain은 fusion하되, register pressure threshold를 넘기면 split한다.

---

예시 2. Row-wise Reduction

관찰:

row length가 작을 때는 warp shuffle reduction이 빠르다.
row length가 커지면 block-level partial reduction이 필요하다.

해석:

reduction 범위가 warp boundary를 넘는 순간 communication structure가 바뀐다.

codegen rule:

row length 기준으로 one-warp-per-row와 one-block-per-row variant를 선택한다.

---

예시 3. Memory-heavy Kernel

관찰:

global memory dependent warp는 ready warp supply가 부족할 때 progress가 크게 감소한다.

해석:

memory latency는 단일 warp 내부에서 해결되지 않고, 다른 ready warp 또는 독립 instruction supply로 숨겨져야 한다.

codegen rule:

memory-heavy kernel은 충분한 active warp 수를 확보하고, load/compute interleaving을 적용한다.

---

예시 4. Shared Memory Bank Conflict

관찰:

특정 stride에서 shared memory latency spike가 발생한다.

해석:

shared memory layout이 bank mapping과 충돌한다.

codegen rule:

shared tile layout에 padding을 삽입하거나 access pattern을 변경한다.

---

11. 최소 실험 범위

초기 계획을 유지하기 위해 GPU probing의 1차 범위는 다음 정도로 제한하는 것이 좋다.

1. Warp execution signature
2. Warp memory access signature
3. Warp-level reduction / communication
4. Block 내부 warp composition effect
5. Shared memory 사용에 따른 trade-off

이 범위는 AI compiler의 kernel generation decision과 직접 연결된다.

반대로 다음 영역은 1차 범위에서는 제외하는 것이 좋다.

SM 간 scheduling 차이
L2 cache slice별 차이
memory controller contention
grid-level block scheduling
multi-kernel interference
runtime context scheduling
PCIe / CPU-GPU transfer 영향

이 영역들은 흥미롭지만, 초기 codegen rule을 만드는 데에는 너무 넓고 간접적이다.

---

12. 문서화 기준

앞으로 각 실험은 다음 질문을 통과해야 한다.

이 실험 결과가 실제 kernel code generation rule로 바뀔 수 있는가?

바뀔 수 있다면 실험한다.

바뀌기 어렵다면 1차 범위에서는 미룬다.

각 실험 문서는 다음 구조를 갖는 것이 좋다.

1. Experiment Purpose
2. Controlled Variables
3. Measured Values
4. Observed Signature
5. Bottleneck Interpretation
6. Codegen Decision Impact
7. Candidate Rule
8. Limitation

예시:

Experiment Purpose:
Elementwise fusion depth가 throughput과 register pressure에 주는 영향을 측정한다.

Observed Signature:
Fusion depth 4까지 throughput이 증가하고, 16부터 감소한다.

Bottleneck Interpretation:
초기에는 global memory traffic 감소가 지배적이고, 이후 register pressure와 occupancy 감소가 지배적이다.

Codegen Decision Impact:
Elementwise chain fusion depth를 무제한으로 늘리지 않고 threshold 기반으로 제한한다.

Candidate Rule:
Register count 또는 occupancy estimate가 threshold를 넘으면 kernel을 split한다.

---

13. 최종 정리

GPU probing은 수치 자체를 얻기 위한 실험이 아니다.

목표는 수치를 통해 다음 연결고리를 만드는 것이다.

operation pattern
→ execution signature
→ bottleneck classification
→ cost model
→ kernel variant selection
→ generated CUDA code

따라서 이 프로젝트에서 probing 결과는 다음 방식으로 사용되어야 한다.

연산 A가 연산 B보다 몇 % 빠르다

가 아니라,

이 연산 구조는 어떤 병목을 만들고,
그 병목을 줄이기 위해 어떤 코드 생성 전략을 선택해야 하는가

로 환원되어야 한다.

한 줄로 정리하면 다음과 같다.

GPU probing은 AI compiler가 kernel variant를 선택하기 위한 empirical cost model을 만드는 과정이다.

그리고 최종 산출물은 실험 결과표가 아니라 다음 형태여야 한다.

Probe Result → Codegen Rule

이 구조가 잡히면 GPU probing은 단순한 성능 관찰이 아니라, 실제 AI compiler의 kernel generation 정책을 만드는 근거가 된다.