(cuda) 2. hello CUDA 출력해보기

hello CUDA 출력해보기

새로운 프로젝트를 생성할 때 임의로 작성된 코드들을 모두 지우고, 아래와 같이 hello CUDA를 출력하는 코드를 작성해보자.

hello-cpp.cu

#include <stdio.h>

void hello(void) {
  printf("hello world!\n");
}

int main(void) {
  hello();
  return 0;
}

C++ 프로그래밍을 해본 사람이라면 많이 익숙한 코드일 것이다.

내용은 C/C++ 내용 뿐이지만 컴파일/실행이 잘되는 것을 확인할 수 있다.

---

프로그래밍 팁

CUDA에서 printf() 출력 문제 해결

커널 안에서 호출한 printf()는 바로 화면에 출력되지 않고 GPU 쪽 버퍼에 저장된다. 이 버퍼는 커널 실행이 끝나고 호스트와 동기화되는 시점에 비로소 화면에 출력되는데, 커널 실행은 비동기라서 동기화 없이 프로그램이 먼저 끝나버리면 출력이 아예 나오지 않을 수 있다. 호스트 쪽 fflush(stdout)는 GPU의 printf 버퍼를 비우지 못하므로, 커널 실행 후 cudaDeviceSynchronize() 같은 동기화 호출을 넣어야 한다.

hello <<<1,1>>> ();
cudaDeviceSynchronize();

이렇게 하면 커널이 끝날 때까지 기다렸다가 printf 버퍼의 내용이 화면에 출력된다.

---

이제 GPU에서 “hello CUDA!” 메시지를 화면에 출력하는 간단한 CUDA 프로그램 작성해보자.

hello-cuda.cu

#include <stdio.h>
#include <cuda.h>

__global__ void hello(void) {
	printf("hello CUDA!\n");
}

int main(void)
{
	hello <<<1,1>>> ();
	cudaDeviceSynchronize();
	return 0;
}

• __global__: 이 키워드는 함수가 GPU에서 실행되는 CUDA 커널임을 나타낸다.
• hello <<<1,1>>> ();: hello 커널을 GPU에서 호출하는 부분. <<<1,1>>>는 커널 실행 구성을 나타내는데, 1개의 블록과 블록 당 1개의 스레드로 커널을 실행함을 의미한다. 총 1개(1x1)의 스레드(코어)가 hello 함수를 실행하게 된다.
• 이 코드는 ‘Linux’에서는 실행이 안될 수 있는데, 그 부분은 뒤에서 다루자.

이제 함수를 여러번 ‘동시 실행’하는 간단한 ‘병렬 컴퓨팅’ 코드를 작성해보자.

hello-parallel.cu

#include <stdio.h>

__global__ void hello(void) {
	printf("hello CUDA%d!\n", threadIdx.x);
}

int main(void)
{
	hello <<<1,8>>> ();
	cudaDeviceSynchronize();
	return 0;
}

• __global__: 함수가 GPU에서 실행된다는 표시.
• threadIdx.x: 현재 스레드의 x 인덱스.
• hello <<<1,8>>> ();: GPU에서 hello 커널 실행. 1개 블록, 블록 당 8개 스레드. 8(1x8)개 core.

코드를 실행해보면 0 ~ 7까지 출력되는 것을 볼 수 있는데, 이것은 하나의 쓰레드가 8번 돈 것이 아니라 8개의 쓰레드가 각각 1번씩 동작한 것이다.

이제는 hello «<1,8»> (); 을 hello «<8,2»> ();로 바꿔보자.

• «<8,2»>
  • 8 x 2 = 16번 (16개의 코어 사용)
• 함수를 8세트, 2번씩 동시 실행

실행해보면 0 ~ 1 (2번)까지 8번 출력되는 것을 확인할 수 있다.