5.1 컴파일 파이프라인 개요

고수준 작업을 효율적인 가속기 코드로 변환하려면 정교한 컴파일 파이프라인이 필요합니다. WIA-AI-011은 프런트엔드 파싱, IR 생성, 최적화 패스, 백엔드 코드 생성 및 디바이스별 어셈블리를 포함하는 다단계 컴파일 프로세스를 정의합니다.

// WIA-AI-011 컴파일 파이프라인
┌─────────────┐
│  고수준 코드 │ (Python, C++, WIA DSL)
└──────┬──────┘
       │ 프런트엔드 파싱
       ↓
┌─────────────┐
│  고수준 IR   │ (텐서 연산, 제어 흐름)
└──────┬──────┘
       │ 최적화 패스 (연산자 융합, 상수 전파)
       ↓
┌─────────────┐
│  중간 IR     │ (디바이스 독립적)
└──────┬──────┘
       │ 백엔드 선택 (GPU, TPU, NPU)
       ↓
┌─────────────┐
│  저수준 IR   │ (디바이스별 명령어)
└──────┬──────┘
       │ 코드 생성 & 최적화
       ↓
┌─────────────┐
│ 기계 코드/PTX│
└─────────────┘

5.1.1 프런트엔드: 소스에서 IR로

// WIA DSL 예제
kernel matmul_kernel(
    input A: tensor,
    input B: tensor,
    output C: tensor
) {
    // 고수준 텐서 표기법
    C = A @ B;  // 행렬 곱셈
}

// 컴파일러가 생성하는 고수준 IR
func @matmul_kernel(
    %A: tensor,
    %B: tensor
) -> tensor {
    %C = tensor.matmul %A, %B : (tensor, tensor)
                                -> tensor
    return %C
}

5.2 중간 표현 (IR) 설계

잘 설계된 IR은 하드웨어 세부 정보를 추상화하면서 이식 가능한 최적화를 가능하게 합니다. WIA-AI-011은 다중 레벨 IR 접근 방식을 채택합니다.

5.2.1 SSA 형식과 제어 흐름

// SSA (Static Single Assignment) 형식
func @forward_pass(%input: tensor) -> tensor {
    // 각 값은 정확히 한 번만 할당
    %0 = conv2d(%input, %weight1)       // 컨볼루션
    %1 = batch_norm(%0, %bn1_params)    // 배치 정규화
    %2 = relu(%1)                       // 활성화
    %3 = max_pool(%2, kernel=2)         // 풀링
    %4 = conv2d(%3, %weight2)           // 두 번째 레이어
    %5 = batch_norm(%4, %bn2_params)
    %6 = relu(%5)
    %7 = flatten(%6)                    // 평탄화
    %8 = matmul(%7, %fc_weight)         // 완전 연결
    %9 = softmax(%8)                    // 출력
    return %9
}

// 조건부 실행
func @conditional_compute(%x: tensor, %threshold: f32) {
    %norm = reduce_sum(abs(%x))

    // 제어 흐름
    cond_br %norm > %threshold, ^compute, ^skip

^compute:
    %result = expensive_operation(%x)
    br ^exit(%result)

^skip:
    br ^exit(%x)

^exit(%final: tensor):
    return %final
}

5.3 연산자 융합 (Operator Fusion)

여러 작업을 단일 커널로 융합하면 메모리 트래픽과 시작 오버헤드가 줄어듭니다.

// 융합 전: 3개의 별도 커널
%0 = conv2d(%input, %weight)    // 커널 1: 컨볼루션
%1 = add(%0, %bias)              // 커널 2: 바이어스 추가
%2 = relu(%1)                    // 커널 3: 활성화

// 문제점:
// - 3번의 글로벌 메모리 쓰기/읽기
// - 3번의 커널 시작 오버헤드

// 융합 후: 1개의 커널
%result = fused_conv_bias_relu(%input, %weight, %bias)

// 이점:
// - 1번의 글로벌 메모리 쓰기
// - 1번의 커널 시작
// - 중간 결과가 레지스터에 유지

// 융합 패턴 인식
pattern ConvBiasRelu {
    match:
        %conv = conv2d(%input, %weight)
        %bias_add = add(%conv, %bias)
        %output = relu(%bias_add)

    rewrite:
        %output = fused_conv_bias_relu(%input, %weight, %bias)
}

5.3.1 고급 융합 전략

// 수직 융합: 프로듀서-컨슈머 체인
%a = op1(%input)
%b = op2(%a)
%c = op3(%b)
// → fused_op123(%input)

// 수평 융합: 독립적 연산 병합
%x = op_a(%input1)
%y = op_b(%input2)
// → fused_op_ab(%input1, %input2)

// 융합 가능성 분석
bool can_fuse(Operation* producer, Operation* consumer) {
    // 1. 메모리 풋프린트 확인
    if (intermediate_size(producer) > register_capacity()) {
        return false;  // 레지스터에 맞지 않음
    }

    // 2. 데이터 의존성 확인
    if (has_other_consumers(producer)) {
        return false;  // 다른 소비자가 중간 결과 필요
    }

    // 3. 융합 이득 추정
    float benefit = saved_memory_bandwidth(producer, consumer);
    float cost = increased_register_pressure();

    return benefit > cost * 1.5;  // 50% 이상 이득 필요
}

5.4 자동 튜닝 및 커널 선택

최적의 커널 매개변수는 문제 크기, 데이터 레이아웃 및 하드웨어 특성에 따라 다릅니다.

// 행렬 곱셈 커널 구성 공간
struct MatmulConfig {
    int tile_M;        // 16, 32, 64, 128
    int tile_N;        // 16, 32, 64, 128
    int tile_K;        // 8, 16, 32
    int thread_M;      // 4, 8, 16
    int thread_N;      // 4, 8, 16
    int warp_M;        // 2, 4
    int warp_N;        // 2, 4
    bool use_tensor_cores;  // true, false
};

// 자동 튜닝 프로세스
class AutoTuner {
    vector search_space;

    MatmulConfig tune(int M, int N, int K) {
        // 1. 검색 공간 생성
        generate_search_space();

        // 2. 유효하지 않은 구성 제거
        prune_invalid_configs(M, N, K);

        // 3. 성능 모델로 후보 필터링
        auto candidates = filter_by_model(search_space, M, N, K);

        // 4. 상위 K개 구성 벤치마크
        MatmulConfig best_config;
        float best_time = INFINITY;

        for (auto& config : candidates) {
            auto kernel = compile_kernel(config);
            float time = benchmark(kernel, M, N, K, num_runs=100);

            if (time < best_time) {
                best_time = time;
                best_config = config;
            }
        }

        return best_config;
    }

    // 성능 모델 (기계 학습 기반)
    float predict_performance(MatmulConfig& config, int M, int N, int K) {
        // GBT 모델로 성능 예측
        float features[] = {
            M, N, K,
            config.tile_M, config.tile_N, config.tile_K,
            config.thread_M, config.thread_N,
            // ... 더 많은 특성들
        };
        return ml_model.predict(features);
    }
};

5.5 메모리 액세스 최적화

메모리 대역폭은 종종 컴퓨팅 처리량보다 성능을 더 많이 제한합니다.

5.5.1 타일링 (Tiling)

// 타일링 전: 캐시 비효율적
for (int i = 0; i < M; i++) {
    for (int j = 0; j < N; j++) {
        float sum = 0;
        for (int k = 0; k < K; k++) {
            sum += A[i][k] * B[k][j];  // B 접근이 스트라이드됨
        }
        C[i][j] = sum;
    }
}

// 타일링 후: 캐시 친화적
#define TILE 32
for (int ii = 0; ii < M; ii += TILE) {
    for (int jj = 0; jj < N; jj += TILE) {
        // 타일을 캐시로 로드
        for (int i = ii; i < min(ii + TILE, M); i++) {
            for (int j = jj; j < min(jj + TILE, N); j++) {
                float sum = 0;
                for (int kk = 0; kk < K; kk += TILE) {
                    for (int k = kk; k < min(kk + TILE, K); k++) {
                        sum += A[i][k] * B[k][j];
                    }
                }
                C[i][j] = sum;
            }
        }
    }
}

// 3단계 타일링 (레지스터, 공유 메모리, 글로벌 메모리)
__kernel void optimized_matmul() {
    // 1. 글로벌 → 공유 메모리 (타일 크기: 128x128)
    __local float tile_A[128][128];
    __local float tile_B[128][128];

    // 2. 공유 메모리 → 레지스터 (타일 크기: 8x8)
    float reg_A[8];
    float reg_B[8];
    float reg_C[8][8] = {0};  // 누적기

    // 계산
    for (int k_tile = 0; k_tile < K; k_tile += 128) {
        load_tile_to_shared(tile_A, tile_B, k_tile);
        barrier();

        for (int k = 0; k < 128; k += 8) {
            load_tile_to_registers(reg_A, reg_B, tile_A, tile_B, k);

            // 레지스터에서 행렬 곱셈
            for (int i = 0; i < 8; i++) {
                for (int j = 0; j < 8; j++) {
                    reg_C[i][j] += reg_A[i] * reg_B[j];
                }
            }
        }
    }
}

5.5.2 프리페칭 (Prefetching)

// 소프트웨어 프리페칭
__kernel void prefetching_example(__global float* data, int N) {
    int gid = get_global_id(0);

    // 다음 반복의 데이터 미리 로드
    float next_value = data[gid + 32];  // 32워프 앞서 프리페치

    for (int i = gid; i < N; i += get_global_size(0)) {
        float current = next_value;  // 프리페치된 값 사용

        // 다음 값 프리페치
        if (i + get_global_size(0) < N) {
            next_value = data[i + get_global_size(0)];
        }

        // 계산
        float result = compute(current);
        data[i] = result;
    }
}

5.6 명령어 수준 병렬성 (ILP)

현대 가속기는 여러 명령어를 동시에 실행합니다.

// 루프 언롤링
// 언롤링 전
for (int i = 0; i < N; i++) {
    c[i] = a[i] + b[i];
}

// 4배 언롤링
for (int i = 0; i < N; i += 4) {
    c[i+0] = a[i+0] + b[i+0];  // 독립적 - 병렬 실행 가능
    c[i+1] = a[i+1] + b[i+1];
    c[i+2] = a[i+2] + b[i+2];
    c[i+3] = a[i+3] + b[i+3];
}

// 소프트웨어 파이프라이닝
void software_pipelining() {
    // 파이프라인 시작 (프롤로그)
    load(chunk_0);

    for (int i = 1; i < N; i++) {
        load(chunk_i);         // 다음 청크 로드
        compute(chunk_{i-1});  // 이전 청크 계산 (중복!)

        // 로드와 계산이 파이프라인됨
    }

    // 파이프라인 종료 (에필로그)
    compute(chunk_{N-1});
}

// 의존성 체인 분리
// 나쁨: 긴 의존성 체인
float sum = 0;
for (int i = 0; i < N; i++) {
    sum += array[i];  // sum에 의존, 순차적
}

// 좋음: 독립적 누적기
float sum0 = 0, sum1 = 0, sum2 = 0, sum3 = 0;
for (int i = 0; i < N; i += 4) {
    sum0 += array[i+0];  // 4개 누적기가 독립적
    sum1 += array[i+1];  // 병렬 실행 가능
    sum2 += array[i+2];
    sum3 += array[i+3];
}
float sum = sum0 + sum1 + sum2 + sum3;

5.7 양자화 및 혼합 정밀도

낮은 정밀도(INT8, FP16)는 더 빠른 계산과 감소된 메모리 사용을 가능하게 합니다.

// 혼합 정밀도 훈련
void mixed_precision_training() {
    // 1. 모델 가중치를 FP16으로 변환
    half* fp16_weights = convert_to_fp16(fp32_weights);

    // 2. FP16으로 순전파
    half* fp16_activations = forward_pass_fp16(fp16_weights, input);

    // 3. FP32로 손실 계산 (수치 안정성)
    float loss = compute_loss_fp32(fp16_activations, labels);

    // 4. FP16으로 역전파
    half* fp16_gradients = backward_pass_fp16(fp16_activations);

    // 5. FP32 마스터 가중치 업데이트
    for (int i = 0; i < num_params; i++) {
        fp32_weights[i] -= learning_rate * float(fp16_gradients[i]);
        fp16_weights[i] = half(fp32_weights[i]);  // 다시 FP16으로
    }
}

// INT8 양자화 인식 훈련
void quantization_aware_training() {
    // 순전파 중 가짜 양자화 (fake quantization)
    Tensor forward_with_fake_quant(Tensor x, QuantParams params) {
        // 양자화 시뮬레이션
        float scale = params.scale;
        int zero_point = params.zero_point;

        // FP32 → INT8 → FP32 (그래디언트 흐름 유지)
        Tensor quantized = round(x / scale) + zero_point;
        quantized = clamp(quantized, -128, 127);
        Tensor dequantized = (quantized - zero_point) * scale;

        return dequantized;
    }

    // 추론 시 실제 INT8 사용
    int8_t* quantized_weights = quantize_to_int8(trained_weights, params);
}

5.8 백엔드 코드 생성

백엔드는 최적화된 IR을 디바이스별 어셈블리 또는 기계 코드로 변환합니다.

// GPU 백엔드 (SIMT 모델)
// IR:
func @vector_add(%a: tensor<1024, f32>, %b: tensor<1024, f32>)
    -> tensor<1024, f32>

// 생성된 PTX (NVIDIA GPU):
.entry vector_add_kernel (
    .param .u64 ptr_a,
    .param .u64 ptr_b,
    .param .u64 ptr_c
) {
    // 스레드 인덱스 계산
    mov.u32 %tid, %tid.x;
    mov.u32 %ntid, %ntid.x;
    mov.u32 %ctaid, %ctaid.x;
    mad.lo.u32 %r0, %ctaid, %ntid, %tid;

    // 데이터 로드
    ld.global.f32 %f0, [ptr_a + %r0*4];
    ld.global.f32 %f1, [ptr_b + %r0*4];

    // 계산
    add.f32 %f2, %f0, %f1;

    // 저장
    st.global.f32 [ptr_c + %r0*4], %f2;
    ret;
}

// TPU 백엔드 (시스톨릭 어레이)
// 행렬 곱셈을 시스톨릭 어레이 명령어로 변환
IR: C = matmul(A, B)

TPU Assembly:
    // 가중치를 어레이로 로드
    load_weights A, weight_buffer

    // 입력 스트림
    stream_input B, input_buffer

    // 시스톨릭 실행 (128x128 어레이)
    systolic_execute 128x128

    // 결과 저장
    store_output C, output_buffer

5.9 JIT 컴파일 및 캐싱

JIT(Just-In-Time) 컴파일은 실제 입력 형상 및 속성을 기반으로 런타임 특수화를 가능하게 합니다.

// JIT 컴파일 캐시
class JITCache {
    unordered_map cache;

    CompiledKernel* get_or_compile(
        Operation* op,
        vector input_shapes,
        vector dtypes
    ) {
        // 1. 캐시 키 생성
        string key = generate_cache_key(op, input_shapes, dtypes);

        // 2. 캐시 조회
        if (cache.find(key) != cache.end()) {
            return cache[key];  // 캐시 적중
        }

        // 3. 컴파일 (캐시 미스)
        auto start = high_resolution_clock::now();
        CompiledKernel* kernel = compile_kernel(op, input_shapes, dtypes);
        auto compile_time = duration_cast(
            high_resolution_clock::now() - start
        ).count();

        printf("JIT 컴파일: %s (%.1f ms)\n", key.c_str(), compile_time);

        // 4. 캐시에 저장
        cache[key] = kernel;

        // 5. 디스크에 영구 저장
        save_to_disk(key, kernel);

        return kernel;
    }

    string generate_cache_key(Operation* op, ...) {
        // SHA256(operation_type + shapes + dtypes + compiler_flags)
        return sha256(serialize(op, input_shapes, dtypes, flags));
    }
};

5.10 성능 프로파일링 및 분석

최적화 노력이 집중되어야 할 곳을 이해하려면 상세한 프로파일링이 필요합니다.

// 루프라인 분석
struct LooplineMetrics {
    float arithmetic_intensity;  // FLOP / Byte
    float peak_flops;            // 이론적 최대 FLOPS
    float peak_bandwidth;        // 이론적 최대 대역폭
    float measured_flops;        // 실제 FLOPS
    float measured_bandwidth;    // 실제 대역폭
};

void roofline_analysis(Kernel* kernel) {
    auto metrics = profile_kernel(kernel);

    // 산술 강도 계산
    metrics.arithmetic_intensity =
        metrics.total_flops / metrics.total_bytes;

    // 병목 현상 식별
    float compute_roof = metrics.peak_flops;
    float memory_roof = metrics.peak_bandwidth * metrics.arithmetic_intensity;

    if (compute_roof < memory_roof) {
        printf("컴퓨팅 제한 (%.1f%% 피크 활용)\n",
               metrics.measured_flops / metrics.peak_flops * 100);
    } else {
        printf("메모리 제한 (%.1f%% 대역폭 활용)\n",
               metrics.measured_bandwidth / metrics.peak_bandwidth * 100);
    }
}

// 커널 프로파일링
void profile_kernel_detailed(Kernel* kernel) {
    ProfileData data = {0};

    // 하드웨어 카운터 설정
    enable_counter("instructions_executed");
    enable_counter("memory_transactions");
    enable_counter("warp_divergence");
    enable_counter("shared_memory_bank_conflicts");

    // 커널 실행
    launch_kernel(kernel);

    // 메트릭 수집
    data.instructions = read_counter("instructions_executed");
    data.memory_txns = read_counter("memory_transactions");
    data.divergence = read_counter("warp_divergence");
    data.bank_conflicts = read_counter("shared_memory_bank_conflicts");

    // 분석 및 제안
    if (data.divergence > 0.1 * data.instructions) {
        printf("경고: 높은 워프 분기 (%.1f%%)\n",
               data.divergence / data.instructions * 100);
        printf("제안: 조건문 최소화 또는 마스크 연산 사용\n");
    }

    if (data.bank_conflicts > 0) {
        printf("경고: 공유 메모리 뱅크 충돌 감지\n");
        printf("제안: 패딩 추가 또는 액세스 패턴 재구성\n");
    }
}