GPU 프로그래밍 - CUDA 기초 문법

CUDA Programming 문법

1. 개요

NVIDIA GPU를 이용하여 프로그래밍을 하려면 기본적으로 CUDA 프로그래밍에 대해서 알아야한다.
프로그래밍을 하기 위해서는 프로세서 구조를 어느정도 알아야하는데, 이전까지 포스팅은 GPU의 구조에 대한 설명이었다.
이번 포스팅은 CUDA 프로그래밍에 대한 기본적인 내용을 알아보겠다.

2. 주요 단어들

CUDA 프로그래밍간에 사용하는 단어는 일반 CPU 프로그래밍과 살짝 다르다.
이 부분에 대해서는 미리 좀 알아두면 좋다.

1) Host and Device

Host란 기본적으로 CPU를 말하며, Device는 GPU를 말한다.
이는 메모리를 이야기할 때도 마찬가지인데 Host Memory는 메인보드에 연결된 RAM을 말하고, Device memory는 GPU에 달린 VRAM을 말한다.

2) Kernel

OS 제일 아래에서 돌아가는 그 Kernel이 아니다.
GPU에서 돌아가는 프로그램을 Kernel이라고 하는데, 아마 GPU와 가장 맞 닿아있는 프로그램이라서 그런게 아닌가 싶다.

3. CUDA C++ Language Extensions

원래 C++에는 포함되어있지 않으나 CUDA용 컴파일러인 nvcc에서 사용하는 식별자를 말한다.
기존 C++ 코드들은 원래 C++ 컴파일러를 통해 처리되지만 이 확장들은 NVCC가 처리하게된다.

1) 확장 식별자

• __global__
  Device에서 돌아가는 프로그램, 즉 Kernel 이라는 뜻이다.
  Host에서 호출할수 있는 함수이다. void 타입의 리턴 값을 가지며 class의 멤버가 될 수 없다. 기본적으로 비동기 함수이므로 CPU와 동기화하고 싶다면 별도의 명령어가 필요하다.

• __device__
  Device에서 돌아가는 프로그램이나, Device에서만 호출이 가능하다. __global__과 __device__는 병행해서 사용할 수 없다.

• __host__
  host에서 돌아가는 프로그램이다, host에서만 호출 할 수 있다. __device__ 와 __host__ 값을 같이 사용할 수 있다.

2) 차원 및 인덱스에 대한 내장 변수

a. 설명

아래 변수는 device에서 구동되는 코드에서만 유효하다.

• gridDim grid가 몇 개의 Threadblock으로 이루어져있는지에 대한 값이다.
  몇 차원으로 구성했는지에 따라 x,y,z로 호출 가능하다.

• blockDim block이 몇 개의 thread로 이루어져있는지에 대한 값이다.
  몇 차원으로 구성했는지에 따라 x,y,z로 호출 가능하다.

• blockIdx
  grid에서 해당 block의 Index를 반환한다.
  몇 차원으로 구성했는지에 따라 x,y,z로 호출 가능하다.

• threadIdx block에서 해당 thread의 Index를 반환한다.
  몇 차원으로 구성했는지에 따라 x,y,z로 호출 가능하다.

• warpSize
  고정적으로 32를 반환한다.

b. 예시

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

// device에서 구동되는 함수
__global__ void kernel() {
    // 전역에서 현재 Thread가 몇번째인지
    int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x;

    // 블록에서 몇번째 Warp인지
    int warpIdInBlock = threadIdx.x / warpSize;

    // 각각의 값에 대해서 출력
    printf("GridDim.x: %d, BlockDim.x: %d, BlockIdx.x: %d, ThreadIdx.x: %d, GlobalThreadId: %d, WarpSize: %d, WarpIdInBlock: %d\n",
           gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, globalThreadId, warpSize, warpIdInBlock);
}

int main() {
    // 블록당 32개의 스레드, 2개의 블록
    int threadsPerBlock = 32;
    int numBlocks = 2;

    // GPU 커널 실행
    kernel<<<numBlocks, threadsPerBlock>>>();

    // 디바이스에서 작업이 끝날때까지 대기 시킴
    cudaDeviceSynchronize();

    return 0;
}

※ 내용 업데이트 및 추가적인 검증 예정

참고문헌

• 서강대학교 임인성 교수님 - 기초 GPU 프로그래밍 수업 자료
• NVIDIA - NVIDIA ADA GPU ARCHITECTURE