[Multi-Core & GPU Programming] CUDA - Stream

"CUDA - Steam"

[Pinned Memory]

▣ CPU-GPU 구조
· GPU Kernel이 아무리 빨라도 데이터를 CPU에서 GPU로 옮기고, 다시 GPU에서 CPU로 옮기는 시간↑

▣ Memcpy Time
> 8192 x 8192 Transpose
· Host Memory → Device Memory 복사 : 44ms
· Kernel 실행 : 2.3ms
· Device Memory → Host Memory 복사 : 44ms
→ Memory 복사 시간 > Kernel 실행 시간

> 2048 x 2048 Matmul
· Host Memory → Device Memory 복사 : 5ms
· Kernel 실행 : 25ms
· Device Memory → Host Memory 복사 : 2.5ms
→ Memory 복사 시간 < Kernel 실행 시간

▣ Memcpy
· CPU가 직접 모든 데이터를 GPU로 복사 X
· DMA(Direct Memory Access)가 Host Memory와 Device Memory 사이 데이터를 전송

** 현실에서는 Host Memory의 일반 Memory(Pageable Memory)에서 바로 DMA를 수행하기 어려움
∵ OS가 Memory Page를 Disk로 Swap Out 할 수 있는데, DMA는 CPU처럼 Page Table을 자유롭게 해석하면서 추적 불가
∴ CPU가 먼저 데이터를 OS가 물리 Memory에 고정해 둔 영역인 Pinned Memory Region으로 복사
→ DMA가 Pinned memory에서 GPU Device Memory로 데이터를 복사

** Pinned Memory가 항상 좋은 것은 아니기 때문에 처음부터 모든 데이터를 Pinned Memory에 두면 안됨
· Pinned Memory는 OS가 해당 Memory를 Disk로 내보내지 못하게 고정
→ 너무 많이 사용하면 OS의 Memory 관리가 어려워지고, 전체 시스템 성능 저하
∴ 큰 데이터 전송이 자주 필요한 경우에만 선택적으로 사용하는 것이 좋음

▣ Pinned Memory 사용
> cudaMallocHost(void** ptr, size_t size)
· 처음부터 Pinned Host Memory를 할당

> cudaHostRegister(void* ptr, size_t size, unsigned int flags)
· 일반 Malloc으로 할당한 Memory를 나중에 CUDA에 등록해서 Pinned Memory처럼 사용

→ 새로 담당하는 경우 : "cudaMallocHost" 사용 / 이미 있는 Host Pointer를 등록하는 경우 : "cudaHost Register" 사용

▶ Memcpy Time With Pinned Memory
> 8192 x 8192 Transpose
· Host Memory → Device Memory 복사 : 44ms ▶ 20ms
· Kernel 실행 : 2.3ms
· Device Memory → Host Memory 복사 : 44ms ▶ 20ms

> 2048 x 2048 Matmul
· Host Memory → Device Memory 복사 : 5ms ▶ 2.5ms
· Kernel 실행 : 25ms
· Device Memory → Host Memory 복사 : 2.5ms ▶ 1.3ms

> 복사 Bandwidth
· 6.5GB/s ▶ 13GB/s

** Pinned Memory는 Kernel 자체를 빠르게 하는 것이 아니라, Host와 Device 사이 데이터 전송을 빠르게 하는 방법

[Stream]

▣ Pipeline
· Memcpy : 주로 PCI-e/NVLINK와 Copy Engine을 사용
· Kernel : GPU의 SM 연산 자원을 사용
→ Memcpy와 Kernel은 서로 다른 자원 사용하므로 Pipeline 가능

▣ CUDA Stream
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
cudaStream_t stream1, stream2, stream3, stream4;
 
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaStreamCreate(&stream3);
cudaStreamCreate(&stream4);
 
cudaMemcpyAsync(d_M1.elements, M1.elements, matSize, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_M2.elements, M2.elements, matSize, cudaMemcpyHostToDevice, stream2);
cudaMemcpyAsync(d_M3.elements, M3.elements, matSize, cudaMemcpyHostToDevice, stream3);
cudaMemcpyAsync(d_M4.elements, M4.elements, matSize, cudaMemcpyHostToDevice, stream4);
 
MatMulKernel<<<dimGrid, dimBlock, 0, stream1>>>(d_M1.elements, …
MatMulKernel<<<dimGrid, dimBlock, 0, stream2>>>(d_M2.elements, …
MatMulKernel<<<dimGrid, dimBlock, 0, stream3>>>(d_M3.elements, …
MatMulKernel<<<dimGrid, dimBlock, 0, stream4>>>(d_M4.elements, …
 
cudaMemcpyAsync(P1.elements, d_P1.elements, matSize, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(P2.elements, d_P2.elements, matSize, cudaMemcpyDeviceToHost, stream2);
cudaMemcpyAsync(P3.elements, d_P3.elements, matSize, cudaMemcpyDeviceToHost, stream3);
cudaMemcpyAsync(P4.elements, d_P4.elements, matSize, cudaMemcpyDeviceToHost, stream4);
cs

1 : 여러 개의 Stream 생성
3~6 : 초기화
8~11 : 각 Stream에 Host Memory 값을 Device Memory에 복사
    ** CPU가 Memory 복사 완료를 기다리지 않고 바로 다음 코드를 실행할 수 있도록 하여 동시 복사 허용
        → cudaMemcpy는 동기식으로 복사가 끝날 때까지 CPU가 대기
    ** "cudaMemcpyAsync"로 실제 비동기 복사를 하려면 Host Memory가 Pinned Memory여야 함
13~16 : 각 Stream에 Kernel Launch 할당
18~21 : 각 Stream에 Device Memory 값을 Host Memory에 복사

· 서로 다른 Stream의 순서 중요 X / 같은 Stream 내의 내부 순서 중요 O
· 같은 Stream 안의 작업은 직렬로 진행 / 서로 다른 Stream 작업은 병렬로 진행
· Stream을 지정하지 않은 CUDA 호출은 Default Stream에 할당
    → Default Stream을 섞어 쓰면 의도치 않게 동기화가 발생할 수 있으므로 주의 필요

> "cudaStreamSynchronize(stream1);" : Stream 1에 들어간 모든 작업이 끝날 때까지 CPU가 대기
> "cudaDeviceSynchronize( );" : 모든 Stream의 모든 작업이 끝날 때까지 대기
> NVVP(NVIDIA Visual Profiler) : 작성한 Code가 실제 HW에서 어떻게 실행되는지 확인 가능한 NVIDIA의 시각화 Profiling 도구

** Stream을 사용하면 한 데이터의 Kernel 실행 중에 다른 데이터의 복사 작업을 수행 가능
    → GPU 연산 자원과 Copy Engine을 동시에 활용해서 전체 실행 시간 감소 가능

· Kernel : 순서 상관 X(1 → 2 → 3 / 1 → 3 → 2)
· Stream 1에서 D ▶ H에 의한 DMA 사용 → Resource Conflict → Stream 1의 DMA 사용 종료 후 Stream 4의 Copy 진행
    → Stream 4가 먼저 H ▶ D에 의한 DMA 사용이 가능하며 이 경우 Stream 4의 DMA 사용 종료 후 Stream 1의 Copy 진행

[Reference]

· 12_cuda_stream_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...
자세한 내용 보기