Source–SASS 관계 기반 GPU 의미 분석 및 최적화 탐지 시스템

1. 개요

본 프로젝트는 CUDA 커널의 원본 소스 코드가 컴파일 과정을 거쳐 PTX와 SASS로 변환되고, 최종적으로 GPU 하드웨어에서 실행되는 전체 과정을 추적하는 분석 시스템을 구축하는 것을 목표로 한다.

기존의 GPU 분석은 주로 다음과 같은 개별 정보 확인에 집중한다.

• 커널 실행 시간
• 레지스터 및 공유 메모리 사용량
• PTX 및 SASS 명령어
• 메모리 병목
• 파이프라인 활용률
• warp stall 원인
• source line과 assembly 간의 단순 대응

본 프로젝트에서는 이러한 정보들을 독립적으로 확인하는 데서 나아가 다음 관계를 하나의 분석 구조로 연결하려 한다.

CUDA Source
→ Source Expression
→ Mathematical Meaning
→ Parallel Algorithm
→ PTX
→ SASS
→ Register / Memory Dependency
→ Runtime Metrics
→ Optimization Diagnosis

궁극적인 목표는 SASS를 단순한 기계어 목록으로 읽는 것이 아니라, 다음과 같은 질문에 답할 수 있는 분석 체계를 만드는 것이다.

• 이 SASS 구간은 원본 CUDA 코드의 어떤 표현에서 생성되었는가?
• 여러 CUDA 문장이 하나의 SASS 구간으로 융합되었는가?
• 하나의 고수준 연산이 몇 개의 저수준 실행 구간으로 분리되었는가?
• 반복되는 FFMA 명령어는 어떤 수학적 수식을 구현하는가?
• 해당 연산은 dot product, reduction, GEMM 또는 다른 연산 중 무엇에 가까운가?
• 고수준 operator의 경계가 실제 실행에서도 유지되는가?
• 컴파일 결과에 직렬 의존성, 불필요한 메모리 이동 또는 비효율적인 명령어 패턴이 존재하는가?
• 어떤 source expression이 실제 성능 문제와 연결되는가?

이를 통해 본 프로젝트는 원본 코드의 의미와 실제 GPU 실행 구조를 연결하는 의미 기반 GPU 정적·동적 분석 시스템을 지향한다.

---

2. 현재까지의 분석 과정

현재 진행 중인 작업은 NVIDIA가 이미 제공하는 컴파일 옵션, 디버그 메타데이터, 바이너리 분석 도구 및 프로파일러를 조합해 필요한 관계 정보를 추출하는 과정이다.

대표적인 정보 출처는 다음과 같다.

nvcc compile options
- line information
- intermediate file preservation
- ptxas resource information

cuobjdump / nvdisasm
- ELF / CUBIN 추출
- PTX 및 SASS 출력
- instruction address
- function-level machine code

Nsight Compute
- source / PTX / SASS 대응
- instruction-level execution metrics
- memory throughput
- stall reason
- pipeline utilization

현재 단계의 핵심은 새로운 컴파일러를 처음부터 만드는 것이 아니라, 컴파일러와 프로파일러가 이미 생성한 메타데이터를 수집하고 이를 일관된 분석 구조로 재구성하는 것이다.

컴파일 과정에서 관계 정보 보존
→ 바이너리 및 프로파일 결과 추출
→ 명령어 단위 구조화
→ Source–SASS 관계 복원
→ 의미 및 성능 패턴 분석

따라서 현재의 GPU probing 작업은 다음과 같이 정의할 수 있다.

컴파일러와 프로파일러가 생성한 중간 정보와 실행 정보를 수집하여, CUDA source에서 실제 GPU 명령어 실행까지의 관계를 복원하는 과정

---

3. 다음 분석 단계: 명령어에서 수학적 의미로

현재까지의 분석이 주로 다음 질문을 다뤘다면,

어떤 CUDA 소스 라인에서 어떤 SASS 명령어가 생성되었는가?

다음 단계에서는 다음 질문을 다루게 된다.

이 SASS 명령어들의 조합은 어떤 수학적 연산을 수행하는가?

예를 들어 CUDA 코드에 다음 연산이 존재한다고 하자.

acc += a[k] * b[k];

수학적으로는 다음과 같다.

[
acc = acc + a_k b_k
]

반복문 전체를 고려하면 다음과 같은 dot product가 된다.

[
acc = \sum_k a_k b_k
]

SASS에서는 대략 다음과 같은 형태로 나타날 수 있다.

LDG    input A load
LDG    input B load
FFMA   acc = A × B + acc
IADD   loop index or address update
ISETP  loop condition
BRA    loop branch

개별 명령어만 보면 단순한 load, multiply-add, integer add, comparison 및 branch에 불과하다.

하지만 이들이 다음 구조와 결합되면 수학적 의미가 형성된다.

두 개의 입력 load stream
+
반복되는 FFMA
+
loop-carried accumulator dependency
+
K 차원 반복
=
dot product

여러 스레드가 서로 다른 출력 위치에 대해 동일한 dot product를 수행하면 행렬 곱셈으로 확장된다.

[
C_{ij} = \sum_k A_{ik}B_{kj}
]

따라서 고수준 operator의 의미는 특정 명령어 하나에 존재하지 않는다.

Operator Meaning
=
Instruction Sequence
+
Data Dependency
+
Loop Structure
+
Memory Layout
+
Thread Mapping
+
Synchronization

이 관점에서 SASS는 단순 명령어 목록이 아니라, 고수준 연산을 구성하는 실행 문법으로 볼 수 있다.

---

4. 고수준 operator가 형성되는 구조

GPU에서 하나의 operator는 계산 명령어만으로 구성되지 않는다.

예를 들어 GEMM은 다음 요소의 결합으로 구현된다.

Global Memory Load
→ Shared Memory Staging
→ Register Load
→ Multiply-Accumulate
→ Register Accumulation
→ Epilogue
→ Global Memory Store

보다 구체적으로는 다음과 같다.

Global Memory
    ↓ LDG / cp.async

Shared Memory
    ↓ LDS

Registers
    ↓ FFMA / HMMA / MMA

Accumulator Registers
    ↓ FADD / FMNMX / conversion

Global Memory
    ↓ STG

따라서 operator는 계산 그래프만으로 정의할 수 없다.

[

Computation Graph
+
Data Movement Graph
+
Thread Cooperation Graph
]

GPU에서는 실제 성능이 데이터 이동과 스레드 협력 구조에 크게 의존하기 때문에, 수식만 복원해서는 충분하지 않다.

예를 들어 같은 행렬 곱셈 수식이라도 다음 요소에 따라 SASS와 성능이 달라진다.

• 한 스레드가 계산하는 출력 원소 수
• accumulator 개수
• global memory 접근 방식
• shared memory tile 배치
• shared memory bank conflict
• vectorized load 사용 여부
• warp 단위 협력 방식
• tensor core 사용 여부
• software pipelining
• loop unrolling
• instruction scheduling

따라서 분석기는 수학적 의미뿐 아니라, 해당 의미가 실제 하드웨어 위에서 어떤 구조로 실현되었는지를 함께 설명해야 한다.

---

5. PMD와의 유사성

PMD와 같은 정적 소스 분석기는 여러 규칙을 한 번의 실행에서 적용하여 다양한 코드 문제를 탐지한다.

그러나 PMD의 핵심은 단순히 여러 규칙을 동시에 실행한다는 점에만 있지 않다.

PMD는 원본 소스를 문자열 그대로 검사하지 않고, 이를 분석 가능한 구조로 변환한다.

Source Code
→ Parser
→ AST
→ Symbol Resolution
→ Type Information
→ Data-flow Analysis
→ Rule Matching

즉, 서로 다른 형태로 작성된 코드도 AST와 데이터 흐름이라는 공통 분석 구조로 정규화한 뒤 규칙을 적용한다.

본 프로젝트 역시 유사한 구조를 가진다.

CUDA Source
→ PTX / SASS Parsing
→ Instruction IR
→ Register Def-Use
→ Control-flow Graph
→ Memory Access Model
→ Semantic Pattern Matching
→ Optimization Rules

PMD와 본 프로젝트를 대응시키면 다음과 같다.

PMDGPU 의미 분석기

AST node	SASS instruction / instruction group
Symbol table	Register and memory operand map
Data-flow analysis	Register definition-use analysis
Control-flow analysis	Branch, predicate and loop analysis
Code smell	Performance smell
Rule violation	Inefficient machine execution pattern
Source location	Source–PTX–SASS relation
Quick fix	Kernel rewrite or optimization suggestion

본 프로젝트는 PMD의 분석 방식을 GPU 저수준 실행 영역으로 확장한 형태로 볼 수 있다.

다만 분석 대상은 코드 스타일이나 일반적인 취약점이 아니라 다음과 같은 GPU 실행 구조다.

• 직렬 FFMA dependency chain
• 부족한 instruction-level parallelism
• 비효율적인 accumulator 구성
• 불필요한 global memory 왕복
• register spill
• shared memory bank conflict
• 비정렬 또는 비병합 memory access
• 불필요한 address calculation
• 불필요한 branch 또는 predicate
• operator fusion 실패
• 낮은 arithmetic intensity
• 파이프라인 활용 불균형

---

6. Lowering과 분석 가능성

고수준 코드에서는 수식, 변수명, tensor shape, operator 이름과 같은 의미 정보가 명확하게 보인다.

반면 실제 하드웨어에서 발생하는 다음 정보는 고수준 코드만으로 확정하기 어렵다.

• 실제 생성된 명령어
• accumulator 개수
• register dependency
• instruction scheduling
• load 및 store 수
• 메모리 공간별 이동
• compiler fusion 여부
• spill 발생 여부
• loop unrolling 결과
• predicate 변환
• tensor core 사용 여부

예를 들어 소스에 다음 코드가 있더라도,

sum += a[i] * b[i];

실제 컴파일 결과가 하나의 accumulator를 사용하는지, 여러 accumulator를 사용하여 ILP를 확보하는지는 SASS를 확인해야 알 수 있다.

Single accumulator
FFMA R9, R2, R3, R9

Multiple accumulators
FFMA R9,  R2, R3, R9
FFMA R10, R4, R5, R10
FFMA R11, R6, R7, R11

따라서 lowering된 표현에서는 고수준 문법보다 실제로 실현된 구현을 분석할 수 있다.

그러나 이를 다음과 같이 표현하는 것은 부정확하다.

고수준에서는 근본적 최적화가 불가능하고 저수준에서만 가능하다.

고수준에서는 다음과 같은 의미 중심 최적화가 더 유리하다.

• algebraic simplification
• dead operator elimination
• common subexpression elimination
• operator fusion
• shape-based optimization
• redundant tensor removal
• algorithm replacement
• constant propagation

반면 저수준에서는 다음과 같은 구현 중심 최적화가 유리하다.

• register dependency 개선
• instruction-level parallelism 확보
• load/store 감소
• memory coalescing
• bank conflict 제거
• instruction scheduling
• pipeline utilization
• register pressure 조절
• spill 방지

따라서 올바른 구분은 다음과 같다.

고수준 표현
→ 의미 중심 최적화

Lowered 표현
→ 구현 중심 최적화

본 프로젝트가 집중하는 영역은 후자이며, 고수준 operator의 형식이나 경계에 제한되지 않고 실제 실행된 구조를 기준으로 최적화 기회를 탐지하는 것이다.

---

7. Lowering과 Normalization의 차이

Lowering은 표현을 실제 실행에 가까운 형태로 구체화하는 과정이다.

Softmax
→ exp + reduction + division
→ PTX
→ SASS

Normalization은 서로 다른 표현을 공통된 분석 형식으로 변환하는 과정이다.

예를 들어 다음 두 명령어는 사용하는 레지스터가 다르다.

FFMA R9, R2, R3, R9
FFMA R12, R5, R7, R12

하지만 분석용 표현에서는 다음과 같이 정규화할 수 있다.

acc_next = input_a × input_b + acc_previous

이 정규화된 표현을 통해 레지스터 번호나 구체적인 주소가 달라도 동일한 패턴으로 탐지할 수 있다.

분석 흐름은 다음과 같다.

Raw SASS
FFMA R9, R2, R3, R9

Normalized Instruction
dst = fma(src0, src1, dst)

Dependency Relation
R9(next) depends on R9(previous)

Pattern
loop-carried accumulation

Possible Mathematical Meaning
Σ a[k] × b[k]

따라서 본 프로젝트에서는 raw SASS를 직접 규칙과 비교하는 것보다, SASS를 다시 분석용 IR로 변환하는 계층이 필요하다.

---

8. Source와 SASS의 다대다 관계

원본 source와 SASS의 관계는 일반적으로 1:1이 아니다.

하나의 source expression이 여러 SASS 명령어로 분해될 수 있다.

y[i] = max(a[i] * b[i] + bias[i], 0.0f);

Address Calculation
LDG a
LDG b
LDG bias
FFMA
FMNMX
STG

반대로 여러 source expression이 하나의 SASS 구간으로 융합될 수 있다.

tmp = a[i] * b[i];
tmp += bias[i];
y[i] = max(tmp, 0.0f);

FFMA
FMNMX
STG

컴파일러 최적화 과정에서는 다음 현상도 발생한다.

• source line 간 instruction 재배치
• 공통 주소 계산 병합
• 여러 연산의 fusion
• 하나의 연산을 여러 basic block으로 분리
• loop unrolling
• dead code 제거
• constant folding
• predicate conversion
• load instruction의 선행 배치
• source 순서와 SASS 순서 불일치

따라서 source와 SASS를 단순 line number 기반으로 1:1 대응시키면 실제 컴파일 관계를 정확히 표현할 수 없다.

필요한 구조는 다대다 관계 그래프다.

Source Node
   ↕
PTX Node
   ↕
SASS Instruction / Region
   ↕
Semantic Operation
   ↕
Runtime Metric

---

9. Source line이 아닌 Source expression

소스 라인 단위 정보는 기본적인 위치 대응에는 유용하지만, 의미 분석에는 부족할 수 있다.

다음 한 줄에는 여러 연산이 포함되어 있다.

y[i] = relu(dot(a, b) + bias[i]);

이 한 줄은 최소한 다음 의미를 가진다.

dot product
bias addition
activation
output assignment
memory store

따라서 source 측 구조도 가능한 경우 다음과 같이 세분화해야 한다.

Source File
→ Function
→ Statement
→ Expression
→ Operator
→ Variable / Tensor

예를 들면 다음 관계를 구성할 수 있다.

Expression: a[k] * b[k]
→ FFMA input operands

Expression: accumulation
→ loop-carried accumulator register

Expression: + bias[i]
→ epilogue FADD or FFMA operand

Expression: relu(...)
→ FMNMX or predicate-select pattern

Assignment: y[i] =
→ global memory store

분석의 실제 목표는 다음 관계에 가깝다.

Source Expression
↔
Lowered Instruction Group

---

10. 관계 그래프의 Edge 종류

각 node 사이에는 단순히 “대응한다”는 관계만 존재하지 않는다.

다음과 같은 여러 관계 유형을 구분할 필요가 있다.

generated_from
lowered_into
split_into
merged_into
fused_with
depends_on
controls
loads_from
stores_to
implements
contributes_to
scheduled_before
scheduled_after
measured_by

예를 들면 다음과 같이 표현할 수 있다.

Source expression
acc += a[k] * b[k]

lowered_into
SASS region 0x120–0x1F0

implements
dot-product accumulation

depends_on
accumulator R9

loads_from
global memory stream A
global memory stream B

Fusion 관계는 다음과 같이 표현할 수 있다.

Source expression: bias addition
Source expression: ReLU

fused_into
SASS region 0x310–0x340

implements
fused bias + activation epilogue

이 관계 모델을 이용하면 단순 위치 대응을 넘어 컴파일러 변환 자체를 표현할 수 있다.

---

11. 분석 가능한 사실과 추론의 구분

모든 분석 결과가 동일한 확실성을 가지는 것은 아니다.

직접 확인 가능한 사실

다음 정보는 컴파일 결과나 프로파일 결과에서 직접 확인할 수 있다.

• source file과 line number
• SASS instruction address
• opcode
• operand
• register definition-use
• branch target
• predicate
• memory space
• load/store instruction
• instruction count
• register usage
• runtime counter
• execution cycle
• stall metric

의미적으로 추론되는 정보

다음 정보는 패턴과 문맥을 바탕으로 추론해야 한다.

• dot product 여부
• GEMM tile 여부
• reduction 여부
• bias broadcast 여부
• operator boundary
• fusion 여부
• tensor 축 의미
• source 변수와 memory stream의 정확한 관계
• 특정 연산의 수학적 목적

따라서 분석 결과에는 confidence와 evidence를 함께 출력하는 것이 적절하다.

Semantic Pattern
Dot-product accumulation

Confidence
0.92

Evidence
- repeated FFMA instructions
- one loop-carried accumulator
- two independent load streams
- source expression contains multiplication and accumulation
- fixed iteration structure

이 구조는 분석기가 불확실한 결과를 확정적인 사실처럼 출력하는 것을 방지한다.

---

12. 내부 분석 구조

전체 분석 시스템은 다음 계층으로 구성할 수 있다.

12.1 Source IR

file
function
statement
expression
variable
tensor
operator
shape
source range

Source IR은 원본 코드의 고수준 의미를 보존한다.

12.2 PTX IR

PTX instruction
virtual register
basic block
memory space
branch
predicate
source annotation

PTX는 CUDA source와 SASS 사이에서 컴파일러 변환을 이해하는 중간 연결 계층으로 사용한다.

12.3 Machine IR

SASS instruction
physical register
instruction address
opcode
operand
predicate
control-flow edge
register def-use
memory access

Machine IR은 실제 생성된 GPU 실행 구조를 표현한다.

12.4 Semantic IR

Machine IR의 여러 명령어를 의미 단위로 묶는다.

address calculation
global load
shared load
shared store
accumulation
reduction
broadcast
activation
normalization
conversion
output store

12.5 Runtime IR

cycle
stall reason
pipeline utilization
memory throughput
cache hit rate
occupancy
warp activity
instruction execution count

12.6 Relation Graph

각 IR 계층을 다대다 관계로 연결한다.

Source IR
↕
PTX IR
↕
Machine IR
↕
Semantic IR
↕
Runtime IR

각 edge는 다음 정보를 포함할 수 있다.

relation type
confidence
evidence
source range
instruction range
analysis rule

---

13. 규칙 기반 탐지 예시

13.1 직렬 FFMA 의존성

Rule
DEPENDENT_FFMA_CHAIN

Evidence
- repeated FFMA
- destination register reused as source
- one accumulator
- loop-carried dependency

Meaning
serial accumulation

Possible Impact
low instruction-level parallelism
latency-bound execution

Suggestion
use multiple independent accumulators

13.2 독립 accumulator 패턴

Rule
MULTI_ACCUMULATOR_ILP

Evidence
- multiple destination accumulators
- no direct dependency between adjacent FFMA instructions

Meaning
instruction-level parallel accumulation

Possible Benefit
higher latency hiding

13.3 Dot product 추론

Rule
DOT_PRODUCT_PATTERN

Evidence
- two input streams
- repeated multiplication-accumulation
- loop or unrolled reduction dimension
- final scalar or partial sum

Inferred Expression
acc = Σ a[k] × b[k]

13.4 Fused epilogue

Rule
FUSED_BIAS_ACTIVATION

Evidence
- accumulation region
- bias load or scalar addition
- activation instruction
- direct output store
- no intermediate global store

Meaning
GEMM + bias + activation fusion

13.5 불필요한 global memory 왕복

Rule
INTERMEDIATE_GLOBAL_ROUNDTRIP

Evidence
- result stored to global memory
- immediately reloaded by a following operation
- producer and consumer have compatible execution scope

Possible Optimization
kernel fusion or shared/register forwarding

13.6 Shared memory bank conflict 후보

Rule
SHARED_BANK_CONFLICT_CANDIDATE

Evidence
- shared memory access
- lane-dependent addresses
- stride aligned with bank count
- repeated access pattern

Possible Impact
serialized shared memory transactions

---

14. Top-down과 Bottom-up 분석

본 시스템은 양방향 분석을 지원해야 한다.

Top-down 분석

원본 수식에서 실제 하드웨어 구현을 추적한다.

Source expression
→ PTX
→ SASS region
→ register dependency
→ memory access
→ runtime behavior

대표 질문은 다음과 같다.

• 이 CUDA 표현은 어떤 명령어를 생성했는가?
• 컴파일러가 fusion을 수행했는가?
• 이 수식은 몇 개의 load와 store를 발생시켰는가?
• 어떤 accumulator 구조로 구현되었는가?
• 실제 병목은 어디에서 발생했는가?

Bottom-up 분석

SASS 패턴에서 원본 의미를 복원한다.

SASS region
→ dependency pattern
→ semantic operation
→ mathematical expression
→ source expression

대표 질문은 다음과 같다.

• 이 FFMA chain은 어떤 수식을 의미하는가?
• 이 명령어 구간은 reduction인가?
• 이 memory access는 어떤 tensor에 대응하는가?
• 이 SASS region은 여러 source operator가 fusion된 결과인가?
• 이 실행 구조는 어떤 고수준 operator의 일부인가?

이 양방향 분석이 가능해야 단순한 disassembler나 profiler를 넘어선 의미 분석기가 된다.

---

15. 권장 실험 순서

복잡한 operator를 곧바로 분석하기보다 기본 연산부터 단계적으로 패턴을 구축하는 것이 적절하다.

1단계: 단일 산술 연산

y[i] = a[i] + b[i];
y[i] = a[i] * b[i];
y[i] = a[i] * b[i] + c[i];

확인 대상:

LDG
FADD
FMUL
FFMA
STG

2단계: 조건 및 activation

y[i] = max(x[i], 0.0f);

확인 대상:

FMNMX
FSETP
SEL
predicate

3단계: 단일 accumulator

for (...) {
    acc += x[i];
}

확인 대상:

loop-carried dependency
FADD chain
loop control

4단계: Dot product

for (...) {
    acc += a[i] * b[i];
}

확인 대상:

two load streams
FFMA chain
accumulator structure

5단계: 다중 accumulator

acc0 += ...
acc1 += ...
acc2 += ...
acc3 += ...

확인 대상:

independent dependency chains
ILP
final reduction

6단계: Reduction

확인 대상:

warp shuffle
shared memory
barrier
tree reduction
atomic

7단계: GEMM

확인 대상:

thread mapping
tile load
shared staging
FFMA or MMA
multiple accumulators
epilogue
store

8단계: Fused operator

GEMM
+ bias
+ activation

확인 대상:

operator boundary disappearance
intermediate store elimination
epilogue fusion

9단계: 복합 operator

Softmax
LayerNorm
Convolution
Attention

이 단계에서는 연산을 여러 semantic region으로 분해한 뒤 전체 operator 관계를 복원한다.

---

16. 예상 출력 형식

분석 결과는 단순 instruction listing보다 규칙, 의미, 근거 및 성능 영향을 함께 보여주는 형태가 적절하다.

Kernel
probe_dependent_ffma

Source
probe_kernels.cu:120–128

SASS Region
0x0080–0x02C0

Detected Pattern
DEPENDENT_FFMA_CHAIN

Semantic Meaning
Serial dot-product accumulation

Mathematical Form
acc = Σ a[k] × b[k]

Machine Structure
- 32 FFMA instructions
- accumulator register: R9
- one loop-carried dependency chain
- two input operand streams

Runtime Evidence
- measured cycle increase
- dependency-related issue stall candidate

Optimization Candidate
Split the accumulation into multiple independent accumulators
and perform a final reduction

Confidence
0.95

Fusion 관계는 다음과 같이 출력할 수 있다.

Source Expressions
- line 210: bias addition
- line 211: ReLU

Mapped SASS Region
0x0310–0x0340

Relation
many source expressions → one fused instruction region

Detected Meaning
Fused bias and activation epilogue

Evidence
- accumulator result reused directly
- FADD or fused bias operand
- FMNMX activation
- no intermediate global store

---

17. 시스템의 한계

SASS 수준으로 내려갈수록 실제 실행 구조는 명확해지지만 원본 의미 정보는 줄어든다.

예를 들어 반복되는 FFMA는 다음 중 어느 것일 수도 있다.

• dot product
• GEMM
• convolution
• polynomial evaluation
• recurrence
• interpolation
• simulation equation

따라서 SASS만으로 고수준 의미를 항상 확정할 수는 없다.

이를 보완하기 위해 다음 정보를 함께 사용해야 한다.

source expression
source line mapping
PTX
register dependency
memory access pattern
loop structure
thread mapping
runtime metric
known kernel structure

또한 컴파일러 최적화로 인해 source 정보가 부분적으로 사라지거나 부정확해질 수 있다.

따라서 분석 결과는 다음 세 수준으로 구분하는 것이 적절하다.

Confirmed
직접적인 메타데이터 또는 명령어 분석으로 확인됨

Strongly Inferred
여러 증거가 일관되게 특정 의미를 가리킴

Candidate
가능성은 있으나 다른 해석도 존재함

---

18. 프로젝트의 최종 정의

본 프로젝트는 단순히 CUDA source와 SASS를 나란히 출력하는 도구가 아니다.

또한 단순한 SASS parser, disassembler 또는 성능 profiler에 그치지 않는다.

본 프로젝트가 지향하는 시스템은 다음과 같이 정의할 수 있다.

원본 CUDA source를 의미의 기준점으로 유지하면서, PTX와 SASS로 lowering되는 과정에서 형성되는 다대다 관계를 복원하고, 명령어 의존성·메모리 이동·제어 흐름·스레드 배치·runtime metric을 분석하여 수학적 연산과 고수준 operator의 의미를 추론하는 GPU 의미 분석 시스템

보다 간결하게 표현하면 다음과 같다.

Source에서 SASS까지의 관계를 기반으로 저수준 GPU 실행 구조를 수식과 operator 수준으로 해석하고, 고수준 문법이나 기존 operator 경계에 제한되지 않는 최적화 기회를 탐지하는 분석기

PMD와 비교하면 다음과 같다.

PMD
Source structure
→ normalized representation
→ rule matching
→ code diagnosis

GPU Semantic Analyzer
Source / PTX / SASS
→ normalized machine representation
→ semantic and performance rule matching
→ mathematical meaning and optimization diagnosis

핵심적인 차별점은 다음 세 가지다.

1. 원본 source와 SASS의 다대다 관계를 보존한다.
2. 명령어를 수학적 의미와 고수준 operator로 복원한다.
3. 고수준 operator 경계가 아닌 실제 하드웨어 실행 구조를 기준으로 최적화 가능성을 탐지한다.

결국 이 프로젝트는 SASS를 단순한 최종 산출물이 아니라, 컴파일러가 선택한 실제 실행 전략을 드러내는 분석 대상으로 바라본다.

그리고 source, lowering 과정, machine code 및 runtime을 하나의 관계 그래프로 연결함으로써 다음을 동시에 가능하게 하는 것을 목표로 한다.

고수준 의미의 보존
+
저수준 구현의 정확한 관찰
+
실행 결과 기반의 검증
+
기존 틀을 벗어난 최적화 탐지