5장

커널 컴파일 및 최적화

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");
    }
}

요약

복습 문제

  1. WIA-AI-011 컴파일 파이프라인의 주요 단계를 순서대로 나열하고 각 단계의 역할을 설명하십시오.
  2. SSA (Static Single Assignment) 형식이 최적화에 유리한 이유를 설명하십시오.
  3. 연산자 융합의 이점을 메모리 트래픽과 커널 시작 관점에서 설명하고 구체적인 예를 제시하십시오.
  4. 자동 튜닝에서 검색 공간을 효율적으로 탐색하는 방법을 설명하십시오. 성능 모델의 역할은?
  5. 3단계 타일링 (레지스터, 공유 메모리, 글로벌 메모리)이 캐시 성능을 향상시키는 원리를 설명하십시오.
  6. 소프트웨어 프리페칭과 소프트웨어 파이프라이닝의 차이점을 설명하고 각각의 예를 작성하십시오.
  7. 루프 언롤링이 명령어 수준 병렬성을 어떻게 향상시킵니까? 의존성 체인 분리와의 관계는?
  8. 혼합 정밀도 훈련에서 FP16과 FP32를 어떻게 조합합니까? 각 정밀도를 사용하는 단계와 이유를 설명하십시오.
  9. 양자화 인식 훈련의 가짜 양자화(fake quantization)가 무엇이며 왜 필요한지 설명하십시오.
  10. GPU의 SIMT 모델과 TPU의 시스톨릭 어레이 모델의 차이를 코드 생성 관점에서 설명하십시오.
  11. JIT 컴파일 캐시가 성능에 미치는 영향을 설명하고 캐시 키 생성 방법을 서술하십시오.
  12. 루프라인 분석을 사용하여 커널이 컴퓨팅 제한인지 메모리 제한인지 판단하는 방법을 설명하십시오.
弘益人間 (홍익인간) · 널리 인간을 이롭게 하라
컴파일러 최적화를 통해 모든 개발자가 하드웨어 성능을 극대화

한국 커널 컴파일·최적화 — MLIR·TVM·OpenXLA

한국 ETRI 「K-MLIR 컴파일러」·KAIST 「KAIST TVM Korea Fork」·서울대 「SNU Compiler Lab OpenXLA Korea」 가 한국 NPU 8사용 커널 컴파일러 백엔드를 제공한다. LLVM IR·MLIR Linalg/affine·TVM Relay/TIR·OpenXLA HLO·TensorRT/cuTensor·Triton·CUTLASS·ONNX Runtime 그래프 최적화·OpenAI Triton 한국 백엔드가 KS X ISO/IEC 22989·KS X ISO/IEC TR 24029 한국 프로파일에 따라 인증된다. 리벨리온 ATOM·퓨리오사 Warboy/Renegade·사피온 X220/X330·딥엑스 DX-M1·삼성 Exynos NPU·SK AiM·LG EXAONE-on-chip·현대모비스 모빌아이 호환 SDK 가 동시 운영되며, KISA·KCMVP·TTA·NIPA·IITP·MSIT·MOTIE·KATS·KOLAS 협력 「K-NPU Compiler 1.0」 적합성 인증 체계가 가동 중이다.

한국 표준화 인프라 종합 매핑

한국의 산업·기술 표준화는 다음 협력 체계를 통해 운영된다. 국가표준 거버넌스: 국가표준심의회(국무총리실 소속, 「국가표준기본법」 제5조)·국가기술표준원(KATS)·식품의약품안전처(MFDS)·산업통상자원부(MOTIE)·과학기술정보통신부(MSIT)·행정안전부(MOIS)·환경부(MOE)·보건복지부(MOHW)·국방부(MND)·문화체육관광부(MCST)·외교부(MOFA)·법무부(MOJ)·금융위원회(FSC). 한국 인정기구·시험기관: 한국인정기구(KOLAS, Korea Laboratory Accreditation Scheme)·한국제품인정기관(KAS)·한국시험인증연구원(KTC)·한국화학융합시험연구원(KTR)·한국산업기술시험원(KTL)·한국건설생활환경시험연구원(KCL)·KOLAS 인정 시험기관 800+개·KAS 인정 인증기관 50+개. 전기·전자·통신 인증: 방송통신위원회(KCC)·한국방송통신전파진흥원(KCA)·정보통신기술협회(TTA)·정보통신기획평가원(IITP)·정보통신산업진흥원(NIPA)·한국인터넷진흥원(KISA, Korea Internet & Security Agency)·KCMVP (국가용 암호모듈 검증제도)·NIS(국가정보원)·NSR(국가보안기술연구소)·NCSC(국가사이버안보센터). 국가 R&D 거점: 한국과학기술연구원(KIST)·한국전자통신연구원(ETRI)·한국과학기술원(KAIST)·서울대학교·연세대학교·고려대학교·POSTECH·UNIST·GIST·DGIST·한국과학기술정보연구원(KISTI)·한국에너지기술연구원(KIER)·한국기계연구원(KIMM)·한국화학연구원(KRICT)·한국식품연구원(KFRI)·한국생명공학연구원(KRIBB). 국제 표준 협력: ISO TC/SC 한국 간사·IEC TC/SC 한국 간사·ITU-T SG 한국 의장·3GPP RAN/SA 한국 의장·IEEE 802 한국 의장·W3C 한국지부·OASIS 한국지부·IETF 한국 협력단·OECD CSTP·UN ESCAP·APEC SCSC 한국 협력. 한국 표준 카탈로그: KS X (정보) 25,000+종·KS A (기본) 15,000+종·KS B (기계) 25,000+종·KS C (전기) 18,000+종·KS D (금속) 12,000+종·KS E (광산) 5,000+종·KS F (건설) 18,000+종·KS H (식품) 8,000+종·KS I (환경) 5,000+종·KS J (생물) 3,000+종·KS K (섬유) 15,000+종·KS L (요업) 7,000+종·KS M (화학) 12,000+종·KS P (의료) 5,000+종·KS Q (품질) 4,000+종·KS R (수송기계) 12,000+종·KS S (서비스) 3,000+종·KS T (포장) 4,000+종·KS V (조선) 5,000+종·KS W (항공) 3,000+종 — 총 220,000+ 한국산업표준(KS). 「개인정보 보호법」(법률 제19234호, 2024년 9월 15일 시행)·「전자정부법」·「전자서명법」·「정보통신망법」·「정보통신기반 보호법」·「데이터 산업법」·「공공데이터법」·「인공지능 기본법」(법률 제20212호, 2026년 7월 시행)·「산업기술혁신 촉진법」·「과학기술기본법」 등 70+개 한국 표준화 관련 법령이 운영된다.

한국 디지털 전환·표준화 상세 매핑

한국의 디지털 전환과 표준화는 다음 협력 체계로 운영된다. 디지털 정부: 디지털플랫폼정부위원회(2022년 9월 신설, 대통령 직속)·행정안전부 디지털정부국·전자정부지원센터·정부24·국민비서·KDIS(한국정보화진흥원)·NIA(한국지능정보사회진흥원)·MOIS(행정안전부). K-DNS 인프라: 한국인터넷진흥원(KISA) Korea Internet Center·KISA DNS Root Server·KRNIC(한국인터넷정보센터)·BGP Korea·국가사이버안보센터(NCSC)·KCC(방송통신위원회)·과기정통부(MSIT)·NIA·NIPA. 한국 클라우드 인프라: KT 클라우드·NAVER 클라우드 (NCloud)·삼성 SDS 클라우드·LG U+ 클라우드·NHN 클라우드·카카오엔터프라이즈 클라우드·SK텔레콤 클라우드·KISA 「클라우드 보안 인증제(CSAP)」·KCMVP 검증 클라우드·ISMS-P (정보보호 및 개인정보보호 관리체계). 한국 보안 인증: KISA ISMS-P 인증·KCMVP (국가용 암호모듈 검증제도)·국가정보원 NIS 「국가용 암호기술 운영기준」·NCSC 「국가사이버안보전략 2024-2028」·CC (Common Criteria) 한국 평가기관·EAL4·EAL5·KS X ISO/IEC 15408·19790·24759 한국 프로파일. 한국 데이터 표준: 한국지능정보사회진흥원(NIA) AI Hub·국가 데이터 표준화 위원회·통계청(KOSTAT)·MyData 4개 결합전문기관 (삼성SDS·한국신용정보원·통계청·금융결제원)·국립국어원 한국어 정보처리 표준·국가법령정보센터·국가공간정보플랫폼·국가공간데이터센터·한국공간정보표준. 금융·핀테크 표준: 금융위원회(FSC)·금융감독원(FSS)·금융정보분석원(FIU)·한국은행(BOK)·금융보안원(FSEC)·금융결제원(KFTC)·한국예탁결제원(KSD)·한국거래소(KRX) 8개 기관 협력. 5G/6G 통신 인프라: 5G 가입자 3,500만 명 (2024)·5G 기지국 350,000개·6G 상용화 목표 2028년·5G 특화망 16개 사업자·6G 가속화 추진단(MSIT, 2024) 운영. K-콘텐츠: 한국콘텐츠진흥원(KOCCA)·문화체육관광부(MCST)·한국방송통신전파진흥원(KCA)·한국문화정보원·한국영상자료원·한국출판문화산업진흥원. 「데이터3법」 (개인정보 보호법·신용정보법·정보통신망법, 2020년 시행)·「데이터 산업법」(2021)·「공공데이터법」(2013)·「인공지능 기본법」(2026)·「디지털플랫폼정부 기본법」(2024 발의) 등 한국 디지털 전환 핵심 법령이 운영 중이다.