[CUDA 공부/책] 2-1.hello_cuda.cu

<CUDA 기반 GPU 병렬 처리 프로그래밍> 책 읽으며 실습 코드 실행해보고 있음.
(코드는 https://github.com/bluekds/CUDA_Programming/tree/master/Book_BJ 여기에서 확인할 수 있음. 근데 짧은 코드면 나는 손으로 직접 작성하는 걸 선호! 그래야 하나하나 따라가면서 이해하는 느낌.)

---

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

__global__ void helloCUDA(void){
    printf("Hello CUDA from GPU!\n");
}

int main(void){
    cudaSetDevice(1);
    printf("Hello GPU from CPU!\n");
    helloCUDA<<<1, 10>>>();
    // GPU 작업 완료 대기
    cudaDeviceSynchronize();
    cudaDeviceReset(); 
    return 0;
}

가장 간단하게 hello world 찍어야지 역시.
이 책의 예제 코드는 leetcuda예제처럼 .cu 를 library로 만든 후, python에서 그 라이브러리를 갖다 쓰는 형식이 아닌,
그냥 nvcc로 컴파일 하면 바로 돌아간당! (우왓! 간편!)

nvcc -o hello_cuda hello_cuda.cu

그래서 사실, 이 간단한 hello world를 찍는 .cu 코드를 library로 만든 후 python으로 실행시켜보려 하였으나...
우선 진도 먼저 나가자라는 마음으로 그냥 있는 코드 찍어봄.

다만!
예제 코드와는 다르게 GPU 작업 완료 대기 코드를 걸지 않으면 Hello GPU from CPU! 만 나오고 끝나버린다!
잠깐 기다리라귯!!
오늘은 Claude의 도움을 받아 그런 이슈를 해결해보았다.
GPU 작업 완료를 대기해보라고 한다.
음... 기다렸더니 잘 되는군.
Hello CUDA from GPU! 가 총 10번 찍힌다.

아, 그리고, 내가 사용하는 환경에 GPU가 여러 대라 그런지?? cudaSetDevice를 안 해주면 안 돌아가서 (아니면 0번이 꽉 차서 그럴까..? 근데 그러기엔 이 아이가 메모리를 뭐 얼마 쓸 것 같진 않은걸... 근데?! 메모리 꽉 찼다는 에러 메시지를 낸다?! 그래서 다른 디바이스로 세팅해줌) GPU 몇 번 쓸 건지 지정해주고 돌려야 돌아갔다.

---

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

__global__ void helloCUDA(void){
    printf("Block %d, Thread %d: Hello CUDA!\n", 
           blockIdx.x, threadIdx.x);
}

int main(void){
    cudaSetDevice(15);
    printf("=== 1 block, 3 threads ===\n");
    helloCUDA<<<1, 3>>>();
    cudaDeviceSynchronize();
    
    printf("\n=== 2 blocks, 2 threads each ===\n");
    helloCUDA<<<2, 2>>>();
    cudaDeviceSynchronize();
    
    return 0;
}

<<<1, 10>>>에 대해서 10이 쓰레드 10개 보낸다는 건 알았는데 앞의 1은 뭐냐 하니까!
늘 들어왔던 그 Grid-Block-Thead의 구조에 대해 설명해주는 우리 클로드군.
그러면서 다른 방식으로도 찍어볼 수 있다고 나에게 팁을 준다.
그렇다면?! 안 찍어 볼 수 없지!
당장 찍어본다!!
<<<1, 3>>> 이면 당연히 3번 찍힐 것 같은데
<<<2, 2>>> 는 어떻게 찍힐까? (4번이겠지)

=== 1 block, 3 threads ===
Block 0, Thread 0: Hello CUDA!
Block 0, Thread 1: Hello CUDA!
Block 0, Thread 2: Hello CUDA!

=== 2 blocks, 2 threads each ===
Block 0, Thread 0: Hello CUDA!
Block 0, Thread 1: Hello CUDA!
Block 1, Thread 0: Hello CUDA!
Block 1, Thread 1: Hello CUDA!

---

Grid.. 그렇다면 Grid가 궁금해지잖아?
Grid는 한 개만 있는 거냐고 물어보니까 그렇다고 한다.

네, Grid는 하나의 커널 실행당 하나만 존재합니다.

하나의 커널 호출 = 하나의 Grid
│
Grid
├── Block (0,0) Block (1,0) Block (2,0)
├── Block (0,1) Block (1,1) Block (2,1)
└── Block (0,2) Block (1,2) Block (2,2)

하지만?!
Grid는 2D, 3D가 가능하다고?!

// 2D
dim3 gridSize(3, 2);  // 3x2 = 6개 블록
dim3 blockSize(4, 4); // 각 블록은 4x4 = 16개 스레드

helloCUDA<<<gridSize, blockSize>>>();

// 3D
dim3 gridSize(2, 2, 2);   // 2x2x2 = 8개 블록
dim3 blockSize(4, 4, 2);  // 각 블록은 4x4x2 = 32개 스레드

helloCUDA<<<gridSize, blockSize>>>();

3차원으로 돌려봤을 때!!

__global__ void helloCUDA(void){
    printf("Grid: (%d,%d,%d), Block: (%d,%d,%d), Thread: (%d,%d,%d)\n",
           gridDim.x, gridDim.y, gridDim.z,      // Grid 크기
           blockIdx.x, blockIdx.y, blockIdx.z,    // 현재 블록 ID
           threadIdx.x, threadIdx.y, threadIdx.z); // 현재 스레드 ID
}

결과는 과연?!

Grid: (2,2,2), Block: (0,0,1), Thread: (0,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (1,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (2,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (3,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (0,1,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (1,1,0)
...
Grid: (2,2,2), Block: (0,1,0), Thread: (2,3,1)
Grid: (2,2,2), Block: (0,1,0), Thread: (3,3,1)

이렇다고 한다.
Grid는 늘 그대로네. Block과 Thread만 바뀌넹.

여튼 이렇게 아주 간단한 hello cuda 찍어보기 + 응용 쪼끔 해봤다리~