[Multi-Core & GPU Programming] MGP Intro

"MGP Intro"


[Python]

· 중간 계층 개수↑ (∵ CPU가 바로 실행할 수 있는 기계어로 직접 돌아가는 언어 X) → 속도↓
· Interpreter 언어 : 한 줄씩 Code를 읽고 해석하면서 실행(C처럼 미리 강하게 최적화된 Native Machine Code로 Compile X)
· Dynamic Type(= Type 미지정) : 변수 a가 int/float/string 등 무엇인지 미리 모르고 실시간 해석 (C는 int a / float a)
· 자동 메모리 관리 & Garbage Collection → 편하지만 비용 발생
· Interface 역할 : 복잡한 고성능 시스템을 사람이 쉽게 다룰 수 있도록 해주는 Front 언어
  → Python으로 다루지만 실제 핵심 연산은 C/C++/CUDA/Fortran 등이 처리(NumPy, PyTorch, TensorFloa, Pandas 등)

※ PyTorch : Python으로 DL 모델을 쉽게 만들게 해주고, 실제 무거운 계산은 빠른 백엔드(C++/CUDA)가 처리하게 해주는 도구
    (Tensor 연산 Library / 자동 미분 지원 / GPU 가속 지원 / 신경망 모델 구축 도구 제공)

[Java]

· JVM(Java Virtual Machine) 위에서 실행(Program → JVM → OS/CPU) → 속도↓

※ Virtual Machine : Type 안정성, 메모리 관리, 보안 검사, 바운드 체크, 예외 처리, Garbage Collection 등을 담당
· HW와 OS는 제각각 → 만약 프로그램을 완전히 Native로만 배포하면 플랫폼마다 따로 빌드하고 관리해야 함
→ VM이 알아서 해당 기계에 맞게 실행 ∴ 각 HW별로 직접 Targeting(like C/C++)하는 작업↓

[Interchange Loops]

· 반복문의 순서를 바꿔 메모리 접근 패턴을 변경(반복 횟수 자체는 동일)
→ 한 번 가져온 데이터를 빠른 메모리에서 재사용 & Cache 적중률 향상

[Optimization Flags]

· 코드의 순서를 바꾸어 CPU가 쉬지 않게 하기, 반복문 줄이기 등 Single Core의 효율을 향상 시켜 개별 연산 속도↑
→ O3로 최적화를 잘해도, 결국 한 번에 한 줄씩 실행되는 Logic 안에서만 성능을 쥐어 짜냄
→ 즉, CPU Core가 16개 있어도 O3만으로는 나머지 15개 Core는 놀게 됨
∴ 병렬화를 할 수 있도록 Multi Core 최적화 설계 필요

[Core]

· CPU 안에 있는 실제 계산 Engine으로 명령을 실행하는 독립적인 처리 장치
· Single Core : 한 번에 한 작업 흐름만 실행 가능
  → 만약 동시에 2개의 Program 실행 시 1개의 Core가 빠르게 번갈아 실행(동시에 실행되는 것처럼 보임)
· Multi Core : 여러 개의 Program 동시에 실행 가능
  → GPU는 많은 Core가 있기에 반복문을 (0~50, 51~100) 이런 식으로 분할하여 Core에 할당 가능 = 속도↑

[Tiling]

· 거대한 데이터를 한 번에 처리하는 대신, CPU의 Cache 메모리 크기에 딱 맞는 작은 조각으로 나누어 처리
→ 데이터가 너무 크면 CPU가 데이터를 읽어올 때 RAM까지 가야 함(= 속도↓)
→ 데이터를 Cache 크기에 맞도록 쪼개어, 해당 조각에 대한 연산이 완전히 끝날 때까지 Cache 안에서만 데이터를 재사용
∴ Cache 재사용 증가 / 메모리 대역폭 낭비↓ / 메인 메모리 접근 횟수↓

[Compiler Vectorization]

· Scalar 연산 : 숫자 1개와 숫자 1개를 더해 결과 1개 출력 → 4번 더하려면 4번의 명령이 필요
· Vector 연산 : 숫자 4개(or 8개, 16개)를 Vector Register에 담아 단 한 번의 명령으로 동시에 더함

[Nvidia GPU Evolution]

· More Cores : 처리할 데이터가 엄청나게 병렬화되어 대량의 병렬 계산을 수행하기 위함
· Larger Memory Size : 기존 GPU 메모리는 크기가 커진 최신 데이터를 저장 불가 → CPU 메모리에 저장 = 데이터 이동 발생
· New Architecture Feature : 특정 Workload 최적화 HW를 추가

[Parallel Architectures]

· 단순히 좋은 HW를 사용한다고 해서 Program이 자동으로 빨라지는 시대 X
→ Programmer가 Parallel Architectures를 활용하여 성능 향상 필요

<배열의 합 구하기(N개의 숫자 합산)>
· for 구문을 사용해 순차적으로 합산

<작업 분배>
① 전체 데이터 N개를 Processor 개수 P로 나눔
② 각 Processor는 자신에게 할당된 범위의 숫자들만 합산
③ 각 Processor는 자신만의 부분 합 생성
④ 생성된 부분 합 합산

<부분 합 합산 방법 - Simple>
· P0 Processor가 Master가 되어 다른 모든 Processor의 부분 합을 수집하여 합산
→ P0가 합산하는 동안 나머지 Processor들은 놀게 됨

<부분 합 합산 방법 - Efficient>
· 토너먼트 방식처럼 계층적으로 부분 합 합산 = 여러 쌍의 Processor가 동시에 서로의 부분 합을 합산
→ 1개의 Processor에 몰리지 않으므로 Network 병목 현상↓ & 전체 소요 시간↓

[Programming Model / Framework]

· Parallel Programming을 구현하는 도구

▣ CPU 병렬화를 위한 Model
<std::threads>
· 운영체제의 저수준 Thread를 직접 제어
· HW 제어권이 높고 정밀한 Thread 관리 가능
· Thread 생성/동기화/데이터 경쟁 등을 Programmer가 일일이 관리해야 하므로 코드 복잡 가능성↑

<OpenMP(Open Multi-Processing)>
· Compiler 지시어(#pragma omp parallel)를 사용하여 병렬 처리를 구현하는 상위 수준 API
· 기존의 순차적인 Code에 한 두 줄의 지시어만 추가해도 병렬화가 가능하여 생산성↑
· 주로 Multi-Core CPU에서 Loop 단위의 병렬 처리에 강력한 성능 발휘

▣ GPU/NPU 병렬화를 위한 Model
<CUDA(Compute Unified Device Architecture)>
· NVIDIA가 만든 Parallel Computing Platform & Programming Model
· GPU HW의 성능을 극한까지 끌어낼 수 있도록 Memory 계층과 Thread 계층을 직접 제어 가능
· NVIDIA GPU에서만 작동하며, HW 구조에 대한 깊은 이해 필요

<Triton>
· OpenAI에서 개발한 언어로, CUDA보다 추상화 수준이 높으면서도 고성능 GPU Kernel을 작성할 수 있게 해줌
· 복잡한 CUDA Programming 없이도 효율적인 FP16/BF16 연산 등을 구현할 수 있어 최신 LLM 최적화에 많이 사용
· PyTorch와의 호환성이 좋아 AI 연구 및 배포 단계에서 각광받음

[Reference]

· 0_intro_mgp_2026 (MGP) - Yongjun Park

댓글

댓글 쓰기

이 블로그의 인기 게시물

[Mini-NPU RTL] NN Reference Model

이미지
"NN Reference Model" 1. Neuron 단 1개의 Neuron만 존재하며 해당 Neuron이 가중합을 잘 구하고 활성화 함수를 제대로 통과시키는지 확인 → 28x28=784개의 입력 데이터와 784개의 가중치를 각각 곱한 후 ∑, 이후 활성화 함수를 통과시킨 뒤, 1개의 값 출력 → "test_data.txt"에서 데이터를 읽어 Neuron에 직접 넣음 → 가중치와 편향이 "w_1_0.mif", "b_1_0.mif"로부터 정적으로 고정되어 Load neuron_tb.sv) include.v) 공통 상수를 미리 정해두는 파일로, 숫자를 고치면 Network 깊이, Neuron 개수 등을 조절 ReLU.v) 공통 활성화 함수인 ReLU를 하드웨어로 구현한 코드 Sig_ROM.v) 공통 Sigmoid 함수 연산을 실시간으로 계산하지 않고  미리 계산된 결과값 을 메모리에서 꺼내 쓰는  LookUp Table  방식의 구현 Weight_Memory.v) 공통 가중치를 저장하는 메모리 장치로 가중치를 읽고 쓴다. Neuron.v) 공통 입력 데이터를 받아 가중치와 곱하고, 누적 한 후, 편향을 더하고, 활성화 함수를 거쳐 출력하는 연산 Unit 2. Network 수많은 Neuron이 합쳐진 전체 시스템의 데이터 흐름과 제어 Logic이 제대로 작동하는지 확인 → 28x28=784개의 입력 데이터가 30개의 Neuron에게 동시에 전달되면 784x30개의 가중치를 곱해 30개의 결과 출력     (첫 번째 입력이 1 Cycle에 서로 다른 30개의 가중치와 병렬 연산으로 곱해진다. ∴ 784 Cycle) → 30개의 입력이 다음 Layer의 20개 Neuron에 전달되어 30x20개의 가중치를 곱해 20개의 결과 출력 → 20개의 입력이 다음 Layer의 10개 Neuron에 전달되어 20x10개의 가중치를 곱해 10개의 결과 출력 → 10개의 출력은 입력 이미지 0~9의 점수를 나...
자세한 내용 보기

[Mini-NPU RTL] TPU (Study Paper)

이미지
"In-Datacenter Performance Analysis of a Tensor Processing Unit" 1. Introduction to Neural Networks   NN은 입력 값들의 가중 합에 비선형 함수(ex. max(0, value))를 적용하는 단순한 인공 뉴런에 기초한다.   ※ 가중합 : 입력 데이터에 중요도인 가중치 를 곱해서 더한다.   ※ 비선형 함수 : 계산된 합을 그대로 쓰지 않고 변형하는 함수(활성화 함수)   ※ max(0, value) : ReLU(Rectified Linear Unit) 함수이며, 음수면 0, 양수면 해당 값을 그대로 내보낸다.   NN은 계층(Layer) 으로 모이며, 한 계층의 출력은 다음 계층의 입력이 된다.   DNN(Deep Neural Networks) 은 NN에서 계층의 수가 많은 것 을 의미한다.   NN은 개발 단계인 Training(학습) 단계와 실전 단계인 Inference(추론) 단계가 있다.   개발자는 계층의 수 와 NN의 유형 을 선택하고, 학습 과정을 통해 가중치 를 결정한다.   학습이 끝난 후, 추론 단계에서는 소수점 아래 긴 자리까지 필요 없는 경우가 많다.   ∴ Quantization(양자화) 단계에서 단순한 정수( 8-bit )로 바꿔 용량↓ & 속도↑         → 8-bit 정수 곱셈 : 16-bit 부동소수점 곱셈보다 에너지 소모 6배↓ & 칩 면적 6배↓         → 8-bit 정수 덧셈 : 16-bit 부동소수점 덧셈보다 에너지 소모 13배↓ & 칩 면적 38배↓   NN의 유형은 3가지로 볼 수 있다.   ·  MLP(Multi-Layer Perceptrons) : 각 계층은 이전 계층의 모든 출력 값(Fully Conn...
자세한 내용 보기