아래의 참조 자료(https://karl6885.github.io/cuda/2018/11/08/NVIDIA-CUDA-tutorial-1/)에 기본적인 튜토리얼이 모두 나와 있으니 아래 자료 참조

아래 내용은 필자가 공부하고 테스트하면서 알게 된 내용을 keyword 중심으로 정리한 내용이라 다소 맥락이 없게 느껴질 수 있다. 향후 더 공부하고 실제 프로그래밍하면서 더 업데이트될 수 있음

Kernel

CUDA는 C++ 프로젝트와는 독립적으로 돌아가는 코드이며, C++과는 사전에 협의된 —인터페이스된— 방법을 통해 명령과 데이터를 주고 받는다고 이해하면 된다. —CUDA는 C라고 이해하면 쉽다.

그러나 코드 자체는 C++ 프로젝트 내에 작성되기 때문에 일반적인 C++ 코드와 구별해 줄 필요가 있는데, __global__ 이라는 키워드가 그러한 역할을 해준다. 이렇게 __global__ 키워드가 붙은 함수를 일반적으로 커널(kernel)이라 부른다.

// __global__ 키워드가 붙은 함수는 C++ 컴파일러가 아니라 NVIDIA CUDA Compiler가 컴파일한다.
__global__ void CudaHello()
{
  printf("Hello CUDA");
}

// __global__ 키워드가 없는 함수는 C++ 컴파일러가 컴파일한다.
int main()
{
	CudaHello<<<1, 1>>>()
}

C++ 코드에서 CUDA 함수를 호출할 때는 그냥 함수명만 쓰는게 아니라 <<<(blockSize), (threadSize)>>> 을 통해 호출한 함수를 몇 개의 block과 block당 몇 개의 thread로 병렬 처리할 것인지를 알려주어야 한다.

위 코드는 1개의 block에 1개의 thread만 사용한다고 알려줬기 때문에 결과는 다음과 같다.

Hello CUDA

만일 blockSize와 threadSize를 2, 3로 지정했다면, 다음과 같이 총 6(=2x3)개가 출력될 것이다.

Hello CUDA
Hello CUDA
Hello CUDA
Hello CUDA
Hello CUDA
Hello CUDA

Thread Hierarchy

앞선 예제와 같이 동일한 작업을 병렬로 수행하기 위해 CUDA를 쓰는 경우는 없을 것이고, 실제 CUDA를 사용하는 것은 로직은 동일하지만 Input이 다르고 그에 따라 Output이 다른 작업을 병렬로 수행해서 빠르게 결과를 얻기 위함일 것이다.

예컨대 0-9999까지의 값을 갖고 있는 10,000개의 배열이 있고, 해당 배열 원소들의 값을 모두 2배씩 증가시키는 로직을 짠다고 가정하자. 이것을 CPU를 통해 로직을 짜면 아래와 같이 10,000번의 반복을 수행하는 코드가 필요하다 —물론 parallel을 쓸 수도 있지만 일단 논외

for (int i = 0; i < a.length; i++)
{
	a[i] *= 2;
}

CUDA를 사용하면 이 10,000개의 배열에 대해 단 1번에 연산을 마칠할 수 있는데, 이는 각 연산을 10,000개의 Thread에 나눠서 처리할 수 있기 때문이다.

// 10,000개의 Thread가 아래의 연산을 수행하면 결과를 1번에 얻을 수 있다.
a[index?] *= 2;

그런데 이것이 가능하려면 10,000개의 Thread에 대해 배열의 각 index를 부여할 수 있어야 한다. 만일 위의 코드에서 배열의 index 자리에 0과 같은 특정한 상수를 넣는다면 10,000개의 Thread가 모두 배열의 0번 값만 바꿀 것이기 때문이다.

이와 같이 배열의 Index를 부여하기 위해 CUDA는 4가지 형태의 값을 제공하는데 그것이 바로 gridDim, blockIdx, blockDim, threadIdx이다.