[Multi-Core & GPU Programming] Others

"Others"

[TensorCore]

· 행렬 곱셈은 큰 행렬을 한 번에 계산하는 것이 아니라, 작은 곱셈과 덧셈(MAC 연산)의 반복으로 구성

· TensorCore는 작은 행렬 단위의 MAC 연산을 매우 빠르게 처리하기 위한 HW

▣ TensorCore 동작 방식

· Memory 저장 및 입력은 FP16으로 수행 → Memory 사용량과 연산 비용↓

· 곱셈 결과를 FP32 Accumulator에 누적 → 정확도를 어느 정도 유지 가능

→ TensorCore는 낮은 정밀도의 입력을 사용해 빠르게 계산하면서, 누적은 더 높은 정밀도로 수행

▣ TensorCore 특징

· 4x4 행렬 단위의 연산을 빠르게 처리(주로 16-bit 곱셈과 32-bit 누적을 사용)

    → 다만, 모든 GPU가 같은 정밀도 지원 X(GPU 세대와 Model에 따라 지원하는 정밀도 다를 수 있음)

· 직접 Low-Level 명령이나 Warp-Level Macro를 통해 사용 가능

· cuBLAS 또는 cuDNN 같은 Library를 통해 간접적으로 사용 가능

· 실제로 개발자가 TensorCore를 직접 다루는 경우 존재

· 보통 cuBLAS의 GEMM 또는 cuDNN의 Convolution 같은 Library를 사용하면 내부적으로 TesorCore가 활용되는 경우↑

[CUDA Library]

· 필요한 기능이 있을 때마다 매번 CUDA Kernel Code를 직접 작성할 필요 없이, 이미 고도로 최적화된 Library 사용 가능

▣ cuBLAS

· CUDA Basic Linear Algebra Subroutine Library → Vector와 행렬 연산을 위한 CUDA Library

· cublasSgemv : Single-Precision General Matrix-Vector Multiplication

    → Float 행렬 A와 Float Vector x를 곱해서 y를 계산하는 함수 (y = alpha * A * x + beta * y)

· 직접 CUDA Kernel을 작성하지 않아도, cuBLAS를 쓰면 NVIDIA가 최적화한 빠른 행렬/Vector 연산을 사용 가능

▣ cuDNN

· CUDA Deep Neural Network Library → Deep Learning에서 자주 사용하는 연산들을 GPU에서 빠르게 수행하도록 만든 Library

  (Convolution, Pooling, Activation, Normalization 등)

· cudnnConvolutionForward : Neural Network의 Forward Pass에서 Convolution을 수행하는 함수

    → 입력 Tensor(x), Filter(w), Tensor Descriptor, Filter Descriptor 등을 넘기면 cuDNN이 내부적으로 최적화된 Convolution Algorithm을 선택해 실행

· Deep Learning Framework인 PyTorch 또는 TensorFlow도 내부적으로 cuDNN을 사용하는 경우↑

▣ Thrust

· GPU를 위한 STL 같은 Library

    → C++ STL에서 Vector, Sort, Reduce, Transform 같은 기능을 쓰듯이, GPU에서도 비슷한 방식으로 사용 가능하게 함

· thrust::device_vector<int> keys(N) : GPU Device Memory에 int Vector를 생성

· thrust::sort(keys.begin( ), keys.end( )) : GPU에서 정렬 수행

→ 직접 정렬 Kernel을 작성하지 않아도 Thrust가 내부적으로 병렬 정렬 Algorithm을 실행

[Evolution]

▣ Independent Thread Scheduling

· Pre-Volta 구조 : Warp 안의 32개 Thread가 하나의 Program Counter를 공유

    → 한 Warp 안에서는 기본적으로 같은 Instruction을 실행하는 Lockstep 구조

· Volta 구조 : 각 Thread가 자기만의 Program Counter와 Stack을 가질 수 있음

    → Warp 내에서 Thread마다 독립적으로 실행 가능

∴ Warp 안의 Thread들이 완전히 항상 같은 위치에서 실행된다고 가정하면 위험

> Lockstep

· Pre-Volta : Warp가 한쪽 Branch를 실행하고, 그 다음 다른 Branch를 실행한 후, 다시 Reconverge해서 Z를 실행

· Volta : Independent Scheduling 때문에 Thread들이 더 유연하게 진행될 수 있음

    → Warp 내부 Thread들이 반드시 완벽한 Lockstep으로 같은 Instruction을 실행한다고 볼 수 없음

△ Scheduling이 더 유연해지고 Algorithm 설계가 더 자유로워짐

▼ "Warp 안에서는 어차피 동시에 움직이니깐 동기화가 필요 없다"고 생각하고 만든 Pre-Volta 코드가 더이상 안전하지 않음

> __syncwarp( )

· Pre-Volta : Warp 내부 Thread들이 Lockstep으로 움직인다고 가정하고, Warp-Level Reduction에서 동기화 사용 X

· Volta : Thread들이 독립적으로 진행될 수 있으므로, Warp 안에서도 명시적인 동기화 필요

· __syncwarp( ) : Warp 내부 Thread들이 특정 지점에서 서로 동기화되도록 함 (각 단계마다 동기화 사용)

** Shared Memory에 값을 Write, 다른 Thread가 해당 값을 Read → 데이터 공유 → 동기화 중요

▣ Large L2 Cache

· Shared Memory를 사용해서 Global Memory 접근을 줄이는 최적화

· Coalescing을 통해 Memory Access를 정리하는 최적화 

→ 데이터가 이미 L2 Cache에 잘 들어간다면, 성능 향상↓

[Reference]

· 9_others_mgp_2026 (MGP) - Yongjun Park