cuda 프로그래밍 공부 #1

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

#define N 100

// [수정 1] 여러 블록을 쓰기 때문에 전체 좌표(Global Index)를 계산해야 합니다.
__global__ void MatAdd(float *A, float *B, float *C)
{
    // 전체 격자에서의 좌표 계산
    // 식: (현재 블록 번호 * 블록 크기) + 블록 내 스레드 번호
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 열(Col)
    int j = blockIdx.y * blockDim.y + threadIdx.y; // 행(Row)


    /*
    *blockIdx (반 번호): "너 몇반이니?"
    *blockDim (한 반의 정원) : "한 반에 몇 명이 있니?"
    *threadIdx (출석 번호): "그 반에서 몇 번이니?" (예: 5번)
    */

    // 인덱스가 배열 범위를 넘지 않는지 확인 후 연산
    if (i < N && j < N) {
        int index = j * N + i;
        C[index] = A[index] + B[index];
    }
}

int main()
{
    float *A, *B, *C;

    cudaMalloc((void**)&A, N*N*sizeof(float));
    cudaMalloc((void**)&B, N*N*sizeof(float));
    cudaMalloc((void**)&C, N*N*sizeof(float));

    float *a = (float*)malloc(N*N*sizeof(float));
    float *b = (float*)malloc(N*N*sizeof(float));
    float *c = (float*)malloc(N*N*sizeof(float));

    for(int i = 0; i < N * N; i++)
    {
        a[i] = 1.0f;
        b[i] = 2.0f;
    }
    
    // 입력 데이터는 Host -> Device로 보내야 합니다.
    cudaMemcpy(A, a, N * N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(B, b, N * N * sizeof(float), cudaMemcpyHostToDevice);

    // 블록 크기 재설정 (하드웨어 한계인 1024개 이하로)
    // 보통 16x16 (256개) 또는 32x32 (1024개)를 씁니다.
    dim3 threadsPerBlock(16, 16); 

    // 전체 N(100)을 커버하기 위해 필요한 블록 개수 계산
    // (100 + 16 - 1) / 16 = 7.xxxx -> 7개 블록 생성
    dim3 numBlocks(
        (N + threadsPerBlock.x - 1) / threadsPerBlock.x,
        (N + threadsPerBlock.y - 1) / threadsPerBlock.y
    );

    // 커널 실행
    MatAdd<<<numBlocks, threadsPerBlock >>>(A, B, C);

    // [Tip] 커널 실행 에러 체크 (이게 있었으면 왜 안되는지 바로 알 수 있습니다)
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("Kernel Launch Error: %s\n", cudaGetErrorString(err));
    }

    // [수정 5] 결과 데이터는 Device -> Host로 가져와야 합니다.
    cudaMemcpy(c, C, N * N * sizeof(float), cudaMemcpyDeviceToHost);

    // 결과 확인
    printf("Result c[0] = %f\n", c[0]); // 3.000000 나와야 정상

    cudaFree(A); cudaFree(B); cudaFree(C);
    free(a); free(b); free(c);

    return 0;
}

다음과 같은 코드를 작성한다.

파일명은 kernel_sample2.cu로 했다.

nvcc -arch=sm_61 -ccbin g++-8 kernel_sample2.cu -o mat_add2

의 코드를 사용해서 컴파일 한다.

결과는 다음과 같이 나오게 된다.

커널 함수는 식별자 __global__을 사용해 정의한다.

이 코드는 행렬의 덧셈을 실행하는 코드이다.

그리고 코드를 쭉 보다 보면 이상한 부분이 있는데

  MatAdd<<<numBlocks, threadsPerBlock >>>(A, B, C);

왜 처음에는 왜 굳이 템플릿을 이렇게 썼지? 싶었다.

그러나 알아보니 이건 템플릿이 아니다. NVIDIA가 CUDA를 만들면서 C++언어 GPU 실행 설정(Execution Configuration)을 위해 특별히 추가한 확장 문법이였다.

일반적인 C++함수는 func(a,b,c)처럼 데이터만 넘기는데, GPU 커널은 데이터 뿐만 아니라

이 함수를 수행할 일꾼(스레드)를 얼마나 뽑을지에 대한 정보가 반드시 필요하다고 한다.

만약 <<<>>>없이 다 섞어 쓴다면 MatAdd(numBlocks, threadsPerBlock, A, B, C); 이렇게

되면 어디까지가 설정값이고, 어디서부터가 진짜 데이터(A, B, C)인지 구별하기 매우 헷갈린다.

그래서 NVIDIA는 아예 구역을 나눴다.

<<< … >>> 안쪽: “하드웨어 세팅” (몇 명의 일꾼을 쓸 것인가?)

( … ) 안쪽: “함수 파라미터” (어떤 데이터를 처리할 것인가?)

그리고 지금은 두 개만 넣었지만, 사실 이 안에는 최대 4가지 정보가 들어간다.

MatAdd<<< 그리드 개수, 블록크기, (공유메모리 크기), (스트림 번호) >>>(A, B, C);

1번(Dg): 그리드 차원 (필수) – numBlocks

2번(Db): 블록 차원 (필수) – threadsPerBlock

3번(Ns): 동적 공유 메모리 크기 (옵션, 기본값 0)

4번(S) : CUDA 스트림(옵션, 기본값 0)

따라서 커널은 <<<(그리드당 블록 수), (블록당 스레드 수) >>> 라는 특수한 연산자로 호출된다.

이어지는 A, B, C는 디바이스 메모리로 불리며 GPU 쪽 메모리이다.

미리 CUDA 전용 메모리 할당 함수 cudaMalloc으로 메모리를 확보한다.

또한, N * N * 1 은 실행 스레드 수를 나타낸다. 커널 내 threadIdx는 CUDA의 예약 변수이고,

현재 커널이 어느 스레드에서 실행되는지 나타낸다.

결국, 스레드에 할당된 ID가 된다.

---

스레드는 CUDA에서 정의되는 실행의 최소 단위이다. 스레드 자신은 3차원으로 정의된다.

따라서, 스레드 ID는 threadIdx.x와 threadIdx.y 그리고 threadIdx.z 차원을 가진다.

3차원을 정의하는 dim3이라는 전용 구조체가 있으므로, 이를 사용해 스레드 수 등을 정의한다.

dim3 threadsPerBlock(16, 16); //여기서 (x, y) 값을 세팅했다!

dim3는 CUDA에서 정의한 특별한 자료형인데, x,y,z 세 개의 값을 가진다.

threadsPerBlock(16, 16) 이라고 쓰면,

x = 16

y = 16

z = 1 (생략하면 기본값 1) 이렇게 자동으로 세팅되어 저장된다.

MatAdd<<numBlocks, threadsPerBlock>>>(A,B,C);

이 <<< >>> 구문이 실행되는 순간, NVIDIA 드라이버는 threadsPerBlock 안의 x = 16, y = 16 정보를 읽어서 GPU의 모든 스레드에게 뿌려준다.

__global__ void MatAdd(...){
    // blockIdx, threadIdx 등은 선언한 적 없는데 쓰인다.
    // 이건 CUDA가 미리 만들어준 '내장 변수 (Built-in Variable)' 이다.
    int i = threadIdx.x;
}

1. Main : dim3 threadsPerBlock(16,16)
2. Launch: <<<…>>> -> GPU에게 이 설계도를 전달.
3. GPU: 스레드를 생성할 때, 각 스레드에게 자신의 위치표를 붙여줌.
   1번 스레드: “너는 threadIdx.x = 0, threadIdx.y = 0 이야.”
   2번 스레드: “너는 threadIdx.x = 1, threadIdx.y = 0 이야.”
4. Kernel Code: int i = threadIdx.x;
   사용잦는 그저 “내 위치포에 적힌 x값을 i에 넣어라” 라고 코드만 짠 것.
   

---

이러한 스레드를 묶은 그룹을 블록이라고 부른다.

블록도 스레드와 마찬가지로 3차원으로 정의된다.

블록에 포함할 수 있는 스레드 수는 하드웨어마다 정해져 있다.

---

다시 블록을 묶은 것을 그리드라고 부른다.

CUDA 툴킷에서는 기본적으로 CPU 메모리(호스트)와 GPU 메모리(디바이스)를 다른 메모리 공간으로 다룬다.

디바이스와 호스트 양쪽 메모리를 확보하고, 호스트 쪽 메모리 내용을 디바이스 쪽 메모리에 cudaMemcpy 함수로 복사한다.