GPGPU 총정리 - (3)

CUDA Programming Mode 

 

CUDA 는 기본적으로 Host + Device 로 나뉘어진다. 

CPU 를 Host , gpu 쪽을 device 라고 부른다. 프로그램은 항상 host 에서 시작하며 , host 코드가 device 메모리 할당 , 데이터 복사 , kernel launch 같은 일을 담당한다. 

실제로 대량 병렬 계산은device 에서 kernel 이 수행한다. 

 

또 , Host와 Device는 각각 별도의 메모리 공간을 DRAM에 유지한다. 이것이 CUDA 프로그래밍에서 명시적인 데이터 전송(`cudaMemcpy`)이 필요한 근본적인 이유이다.

다만, CUDA는 Unified Memory 기능도 제공한다. Unified Memory를 사용하면 모든 프로세서가 단일 일관된 메모리 이미지(coherent memory image)를 공통 주소 공간에서 볼 수 있도록 managed memory 공간을 정의할 수 있다. 이는 프로그래밍을 단순화하지만, 내부적으로는 여전히 데이터 마이그레이션이 발생한다.

 

 

Host  ( CPU ) 

- CPU 코어들이 c++ 프로그램 실행하면서 CUDA RT API ( runttime api 호출 ) 

- cudaMalloc, cudaMemcpy, cudaFree, 커널 런치 등 모든게 HOST 코드 쪽에서 api 를 통해 실행된다. 

아직 나오지는 않았지만 Pinned memory 영역인데 ,  원래 malloc 으로 할당한 메모리는 pageable 이라 OS 가 swap 할 수 있는데 , pinned memory 는 물리 메모리라서 (locked ) copy Engine 이 DMA 로 직접 전송이 가능하다고 한다 ( 속도 더 빠름 ) 

다만 , 이거는 나중에 설명할 것 같아서 이정도 이야기 하고 생략. 

 

Device 0 

SM 들이 그려져 있다. 안에 물결 모양이 아마도 cuda core 나타내는 것 같고 커널 코드가 SM 위에서 작동하는 형식이다. 

 

연결부 

copy engine : Host memory 와 device 사이의 데이터 전송을 담당한다. ( D2H , H2D 둘다 가능 ) 

이 엔진은 SM 과 독립적으로 작동해서 ,  커널 실행과 데이터 전송을 동시에 수행( OVERLAP ) 할 수 있다. 

 

CONTEXT : 

CUDA context 가 GPU 자원 관리의 첫 단위이다. ( 레지스터 , 메모리 할당 ,커널 등등 .. 상태 관리 ) 

cuda api 호출시 이 컨텍스트 생성에 시간이 걸리는게 GPU initialization overhead의 원인이기도 하다 ( 그래서 과제에 시간 측정을 하는 코드에 warm-up 을 해야하는 것. ) 

 

STREAMS : 

Default : 같은 strema 내 명령은 순서대로 실행된다. 

Non default : 서로다른 stream 의 명령은 동시에 실행될 수 있다.  

 

 

 

Device 1 

GPU 가 2개 이상일 수도 있다는걸 보여주고 싶었던듯 , 

 

대략적으로 이런 방식으로 처리한다고 보면된다.

.cu 파일에는 보통 host code & device code 가 같이 들어가 있다. 

이중에서 기본적으로 알아야할 문법이 

 

• __global__ : GPU에서 실행되는 kernel
• __device__ : GPU에서만 호출 가능한 device 함수
• __host__ : CPU에서 실행되는 함수 (안 쓰면 기본적으로 host 함수) 

사실 우리가 쓰는 c/c++ 에서 함수 앞에 원래는 __host__ 가 붙어있는거다 ! 이정도만 알면된다. 

또, cuda 문법이 들어가 있는 파일은 nvcc 로 컴파일 해야한다. nvcc 는 .cu 파일을 host code / device code 로 분리해서 알아서 잘 컴파일 해준다. 

 

 

Kernel lanch : <<< >>> 문법에 대한 이해 . 

 

vectorAdd<<<grid_dim, block_dim, dynamic_smem_bytes, stream>>>(args); 

보통 이런식으로 자주 쓴다. 
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

 

 

grid_dim : block 이 몇개 있는지 

block_dim : 각 block 안에 thread 가 몇개인지

-------- 위 2개는 필수-------- 

dynamic : block 당 동적 shared memory 크기 

stream : 어떤 CUDA stream 에서 처리할까 

 

 

예를들어서 N 개의 원소 처리 + block 당 256 thread 쓴다고 하면 

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(...);

( 딱 떨어지지 않는 부분 까지 커버하기 떄문 ) 

그리고 kernel 에서는 if (i < N ) 범위로 한번 더 체크하는게 CUDA 에서 거의 기본이라고 한다. 

 

 

일단 먼저 풀 코드를 보고 이해를 해보려고 시도해보자. 

 

 

Sample Programming 

// Vector addition: C = A + B
#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

// ── (1) Host 함수 정의 ──
void vectorAdd_host(const float* A, const float* B, float* C, int numElements) {
    for (int i = 0; i < numElements; i++)
        C[i] = A[i] + B[i];
}

// ── (2) Device 함수(커널) 정의 ── [C++ language extension: __global__]
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;  // [C++ language extension: built-in variables]
    if (i < numElements)
        C[i] = A[i] + B[i];
}

int main(void) {
    int numElements = 8000000;
    size_t size = numElements * sizeof(float);

    // ── (3) Host 메모리 할당 ──
    float *h_A = (float *)malloc(size);
    float *h_B = (float *)malloc(size);
    float *h_C_host = (float *)malloc(size);
    float *h_C_device = (float *)malloc(size);

    // ── (4) 에러 체크 ──
    if (h_A == NULL || h_B == NULL || h_C_host == NULL || h_C_device == NULL) {
        fprintf(stderr, "Failed to allocate host vectors!\n");
        exit(EXIT_FAILURE);
    }

    // ── (5) 데이터 생성 ──
    for (int i = 0; i < numElements; ++i) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }

    // ── (6) Host 함수 호출 (CPU에서 벡터 덧셈) ──
    vectorAdd_host(h_A, h_B, h_C_host, numElements);

    // ── (7) Device 메모리 할당 ── [CUDA Runtime API]
    float *d_A = NULL;
    cudaMalloc((void **)&d_A, size);
    float *d_B = NULL;
    cudaMalloc((void **)&d_B, size);
    float *d_C = NULL;
    cudaMalloc((void **)&d_C, size);

    // ── (8) 데이터 전송: H2D ── [CUDA Runtime API]
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // ── (9) Grid과 Thread Block 설정 ──
    int threadsPerBlock = 256;
    int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;

    // ── (10) 커널 런치 ── [C++ language extension: <<<...>>>]
    vectorAdd <<<blocksPerGrid, threadsPerBlock>>> (d_A, d_B, d_C, numElements);

    // ── (11) 데이터 전송: D2H ── [CUDA Runtime API]
    cudaMemcpy(h_C_device, d_C, size, cudaMemcpyDeviceToHost);

    // ── (12) 결과 검증 ── [C++ language extension: fabs]
    for (int i = 0; i < numElements; ++i) {
        if (fabs(h_C_host[i] - h_C_device[i]) / h_C_host[i] > 1.0e-6) {
            fprintf(stderr, "Result verification failed at element %d!\n", i);
            exit(EXIT_FAILURE);
        }
    }
    fprintf(stdout, "\nTest PASSED\n");

    // ── (13) Device 메모리 해제 ──
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // ── (14) Host 메모리 해제 ──
    free(h_A);
    free(h_B);
    free(h_C_host);
    free(h_C_device);

    fprintf(stdout, "\nDone\n");
    return 0;
}

 

(자세한 참고)  

실행 흐름 상세

위 프로그램의 실행 흐름을 단계별로 추적하면 다음과 같다:

**Step 1 — Host 메모리 할당 (`malloc`)**  

CPU 쪽 메모리에 입력 벡터 `h_A`, `h_B`와 출력 벡터 `h_C_host`(CPU 결과), `h_C_device`(GPU 결과 수신용)를 할당한다. 이 메모리는 일반적인 C의 `malloc()`으로 할당되는 pageable 메모리이다.

**Step 2 — 데이터 생성 (`rand()`)**  

`h_A`와 `h_B`에 0.0~1.0 범위의 랜덤 float 값을 채운다.

**Step 3 — Host 함수 실행**  

CPU에서 순차적으로 벡터 덧셈을 수행하여 `h_C_host`에 저장한다. 이것은 나중에 GPU 결과를 검증하기 위한 참조값이다.

**Step 4 — Device 메모리 할당 (`cudaMalloc`)**  

GPU 메모리(global memory)에 `d_A`, `d_B`, `d_C`를 할당한다. `cudaMalloc`은 CUDA Runtime API 함수로, 첫 번째 인자로 포인터의 주소(`void **`)를 받는다. 이 포인터는 GPU 메모리의 주소를 가리키게 되며, Host 코드에서 직접 역참조(dereference)할 수 없다. 오직 `cudaMemcpy`나 커널 인자로만 사용 가능하다.

```c

float *d_A = NULL;

cudaMalloc((void **)&d_A, size);  // d_A는 이제 GPU 메모리를 가리킴

```

**Step 5 — 데이터 전송: Host → Device (`cudaMemcpy`, H2D)**  

`h_A`와 `h_B`의 데이터를 GPU 메모리의 `d_A`와 `d_B`로 복사한다. `cudaMemcpyHostToDevice` 플래그로 전송 방향을 지정한다. 이 호출은 **동기적(synchronous)**이다 — CPU는 전송이 완료될 때까지 블록된다.

```c

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);  // dst, src, size, direction

cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

```

**Step 6 — Grid/Block 설정 및 커널 런치**  

`threadsPerBlock = 256`으로 각 블록의 스레드 수를 설정하고, `blocksPerGrid`를 계산하여 전체 데이터를 커버하도록 한다. ceiling division 공식 `(N + T - 1) / T`를 사용한다.

```c

int threadsPerBlock = 256;

int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;

// = (8000000 + 255) / 256 = 31250

```

커널 런치는 `<<<blocksPerGrid, threadsPerBlock>>>` 문법으로 수행된다. 이 호출은 **비동기적(asynchronous)**이다 — CPU는 커널을 GPU에 제출하고 즉시 다음 코드로 진행한다.

**Step 7 — 데이터 전송: Device → Host (`cudaMemcpy`, D2H)**  

GPU 결과 `d_C`를 CPU 메모리 `h_C_device`로 복사한다. 이 `cudaMemcpy` 호출은 동기적이므로, GPU에서 커널 실행이 완료되고 데이터 전송까지 끝날 때까지 CPU가 대기한다. 따라서 별도의 `cudaDeviceSynchronize()` 없이도 결과가 안전하게 전달된다.

**Step 8 — 결과 검증**  

CPU 결과와 GPU 결과를 원소별로 비교하여 상대 오차가 1.0e-6 이내인지 확인한다.

**Step 9 — 메모리 해제**  

`cudaFree()`로 GPU 메모리를, `free()`로 CPU 메모리를 해제한다.

 

위 코드를 보면 이제 대략적인 감이 생긴다. 

특히 여기서 코드 짤 때 가장 익숙하지 않았던 부분은 kernel 쪽 코드는 "thread" 기준으로 작성을 해야 한다는 것이다. 

그러니까 기존 c/c++ 에서는 어떤 배열을 관리할때 순서대로 접근하고 , 해당 index 안에서 돌아갈 로직을 결정해주는 것을 명시적으로 코드로 작성해야 했다면 , 

kernel 쪽 함수에서는 thread_id 를 하나 지정하고 , 해당 thread 가 어떤 것을 처리해야 하는지에 대한 관점으로 생각해야한다 ( 추후에 더 설명 ) 

 

 

nvcc 컴파일 과정을 한번 살펴보자. 

 

1. Code separation 

nvcc 가 .cu 파일 받으면 먼저 코드 분리 부터 시작한다. 

판별 기준은 간단하다 . __global__  , __device__ , blockIdx , trheaIdx 같은 CUDA C++ languageExtension 사용하는 부분이 device  / 나머지는 host code 이다. 

 

2. Device code compilation 

두가지정도 소개한다. 

 

PTX :

가상 GPU 아키텍쳐용 어셈블리이다. C++ -> ptx -> sass  정도이다. 

 

SASS : 

실제 gpu 아키텍쳐용 기계어이다. sm_xx 옵션으로 지정할 수 있다. 

 

음.. 이부분은 조금 과한 것 같으니 일단 패스하도록 하겠다. 

 

 

3. Host code compilation 

여기서 nvcc 가 <<< ... >>> 문법을 표준 CUDA Runtime API 호출로 변환한다 

 

예를들어서

// 변환 전 (CUDA extension)
vectorAdd <<<blocksPerGrid, threadsPerBlock>>> (d_A, d_B, d_C, n);

// 변환 후 (표준 C++)
void* args[] = { (void*)&d_A, (void*)&d_B, (void*)&d_C, (void*)&n };
cudaLaunchKernel((const void*)vectorAdd, dim3(blocksPerGrid), 
                  dim3(threadsPerBlock), args, 0, 0);

 

이런식으로 자동으로 변환한다. 

 

변환된 코드를 C++ 컴파일러에 넘긴다. 

 

4. Linking 

host 파일 + device 코드 + cuda runtime library 합쳐서 단일 실행파일을 만든다. 

 

 

정리

[1] .cu 파일 작성
    - host code
    - device code
    - CUDA 문법 포함

[2] nvcc가 분리
    - host 부분 추출
    - device 부분 추출

[3] device 부분 컴파일
    - PTX 생성 (compute_XX)
    - SASS/cubin 생성 (sm_XX)

[4] host 부분 컴파일
    - <<<>>> 를 CUDA runtime API 호출로 변환
    - gcc/clang/cl.exe 로 일반 C++처럼 컴파일

[5] 링크
    - host object + device binary + CUDA runtime
    - 최종 executable 생성 

 

 

 

1. __global__ 

 

실행 위치: device(GPU)
기본 호출 위치: host(CPU)  

CC 5.0 이상에서는 device에서도 호출 가능(동적 병렬성) 

 

__global__ 이 kernel entry function 이다. 

kernel<<<grid, block>>>(args); 

이것 처럼 execution oncifuration 이 필요하다

왜냐하면 , 함수 호출 몇개가 아니라 

" GPU 에 Thread 몇개 , block 몇개를 만들어서 병렬 실행하라 " 라는 요청이기 때문. 

 

쉽게 이야기하면 __global__ 은 단순 함수가 아니라 GPU 작업의 진입점 ( entry point ) 이다.  

 

 

 

4가지 특성 

 

• 반드시 void 반환 — GPU에서 실행되는 함수가 CPU로 값을 직접 return할 방법이 없다. 결과는 device memory에 쓰고, cudaMemcpy로 가져와야 함.
• 실행 구성 필수 — <<<grid, block>>>을 반드시 붙여야 함. "몇 개의 스레드를 어떤 구조로 런치할 것인가"를 지정합니다.
• 비동기 호출 — CPU가 커널을 GPU에 제출하면 즉시 반환합니다. GPU가 아직 실행 중인데 CPU는 이미 다음 줄을 실행해요. 결과를 확인하려면 cudaDeviceSynchronize()나 동기적 cudaMemcpy가 필요합니다. ( 그래서 시간 젤 때 , 명시적으로 cpu 를 돌지 못하도록 막음)  -
  • launch 했다 = GPU에 일 시켰다
  • 끝났다 = 아님
• 재귀 불가 — __global__ 함수는 자기 자신을 호출할 수 없다. (단, CC 5.0+ 에서 Dynamic Parallelism으로 다른 __global__ 함수를 호출하는 건 가능)

 

 

2. __device__ 

GPU 내부 전용 함수이다.

 

• device에서 실행
• device에서만 호출 가능 

예시 ) 
__device__ float helper(float x) { return x * x; }

__global__ void kernel(float *data, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) data[i] = helper(data[i]);  // 커널 안에서만 호출 가능
}

 

 

추가로 device 와 global 을 함께 쓰는 것은 불가능하다. 

 

( __device__ 함수는 kernel이 아니다 -> 그냥 GPU 내부 코드에서 보조 함수처럼 호출한다.  ) 

 

3. __host__ 

우리가 원래 쓰는 c/c++ 함수들이다 ( 그냥 CPU 에서 돌아가는.. ) 

 

__device__  __host__ 같이 쓰면 host 와 device 둘다 컴파일이 된다고 한다. 

근데 주의해야할 점이 CPU / GPU 용 함수가 이렇게 둘 다 생기게 됩니다. 

(nvcc 가 함수를 두번 컴파일 - host / device 용 ) 

 

< 요약 테이블 > 

( 실행 위치 / 호출 가능 위치 ) 

 

 

 

Built-in Vector Types

cuda 가 기본 정수/부동 소수점 묶어서 제공하는 구조체이다. 

GPU 에서 데이터를 2, 3, ,4개씩 묶어서 loag/store 할때 좋다. 

 

예를들어서..

float4 pixel = make_float4(r, g, b, a); // RGBA 색상

float2 coord = make_float2(u, v); //

텍스처 좌표 int2 pos = make_int2(row, col); // 2D 인덱스 

 

대략 이런식으로 사용할 수 있겠다. 

 

그중에서도 dim3 를 가장 자주쓴다. 

 

unit3 기반의 정수 벡터 타입이고 x,y,z 성분을 가진다. 

용도로는 kernel configuration ( grid 와 block 의 차원 지정하고 읽어오는 것. ) - 지정하지 않으면 기본 값은 1로 설정한다. 

cuda 의 grid 와 block 이 1/2/3 D 가 가능하기 때문 . 

 

 

Thread Hierarchy (Prior to CC 9.0) 

여기 파트는 뭐랄까 c 언어 처음배울때로 돌아가서 ( 자료구조처럼 ? ) 배열에 대해서 자리값 계산하고.. 이런 일을 종종 처음에 하는데 , 여기 파트도 그런 느낌이다. 

 

기본 단위에 대한 개념은 많이 다뤘으니 패스하고. 이부분에서는 오로지 블럭 계산을 잘해보기 위함에 목적을 맞춰서 설명해보려고 한다. 

 

 

 

2. built-in variables for dimensions and indices 

built-in 변수 5개  ( device 에서 실행되는 함수안에서만 유효 ) 

변수  | 타입 | 의미

 

gridDim	dim3	grid의 크기 → 블록이 몇 개인지 (x, y, z)
blockDim	dim3	thread block의 크기 → 블록 당 thread 개수
blockIdx	uint3	이 block이 grid 안에서 몇 번째인지
threadIdx	uint3	이 thread가 block 안에서 몇 번째인지
warpSize	int	warp 크기 = 32 (상수)

 

이를 숙지하고 1D / 2D example 들을 확인해보려고 한다. 

 

1D 에서는 비교적 계산이 쉽다. 

 

int x = blockDim.x * blockIdx.x + threadIdx.x;

-> (몇개의 thread 가 한 블럭에 있는지) * (몇번째 블록인지) + ( 해당 thread 안에서 내가 몇번째인지 )

 

이런식으로 나타내면 된다. 처음볼때는 변수명이랑 매칭이 안되서 살짝 헷갈리는데 보다보면 익숙해져서 괜찮다. 

 

 

 

 

 

2D 부터는 살짝 귀찮지만 , 못할정도는 아니다. 

 

총 thread 수 = 4096 x 2160 

block 크기 = 32 x 16 -> ( blockDIM.x x blockDIM.y ) 

block 개수 = (4096 / 32 ) * ( 2160 / 16 ) -> ( gridDim.x * gridDim.y ) 

 

* block 안에서 나는 몇번쨰 trhead 인가 ? 

Thread ID = blockDim.x * threadIdx.y +  threadIdx.x

- x 먼저 채우고 , 다 차면 y 가 올라가는 row-major 순서이므로. 

- 여기서 blockDim x = 32 / y = 16 

- threadIdx 는 해당 블럭에서 x,y 가 지금 어디있는지 알려주므로.. ( 예시에서는 3,2 ) 

- 한 행의 길이(blockDim.x) × 몇 번째 행(threadIdx.y) + 그 행에서 몇 번째(threadIdx.x) 

-> Thread ID = 32 * 2 + 3 으로 계산하면 된다. 

 

 

* 전체 grid 에서 global ID 

x = blockDim.x * blockIdx.x + threadIdx.x ( 열방향 ) 

y = blockDim.y * blockIdx.y + threadIdx.y  ( 행방향 )

 

 

global_id = gridDim.x * blockDim.x * (blockDim.y * blockIdx.y + threadIdx.y) + blockDim.x * blockIdx.x + threadIdx.x 

( Global ID = ROW * gridDim.x * blockDim.x + ROW ) 

 

쉽게 확인해보자. 

x = x축에서의 시작점 + x축에서의 내 로컬 위치
y = y축에서의 시작점 + y축에서의 내 로컬 위치 

(동일)

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y; 

  

-> 여기까지는 Linear index 계산임 ! ( 전역 좌표 계산을 했고 . ) 

이미지가 row-major 로 1차원 배열에 저장되어 있으면 1차원 메모리 index 계산을 해야한다. 

 

다시 정리해보면 

 

---

1) 전역 좌표 계산

이건 그냥 위치를 구하는 거다.

 

x = blockIdx.x * blockDim.x + threadIdx.x;
y = blockIdx.y * blockDim.y + threadIdx.y;

 

여기에는 gridDim.x가 안 들어감. 

 

2) 1차원 메모리 index 계산

이미지가 row-major로 1차원 배열에 저장

 

idx = y * width + x;

 

width ( 한 줄 길이 )

width = gridDim.x * blockDim.x 

 

 ( 먼저 (x, y)를 구함 그 다음 필요하면 idx = y * width + x 로 1차원 인덱스로 바꿈 )

 

 

So,

idx = y * ( gridDim.x * blockDim.x  ) + x; 

 

 

 

(기본 문법 ) 

kernelFunc<<<grid_dim, block_dim, dynamic_smem_bytes, stream>>>(인자들);

• grid_dim: 블록 몇 개 만들지 (필수)
• block_dim: 블록당 스레드 몇 개 (필수)
• dynamic_smem_bytes: 동적 shared memory 크기 (선택, 기본 0)
• stream: 어느 stream에서 실행할지 (선택, 기본 0)

 

<<< >>> -> nvcc 컴파일러가 컴파일하면 API 호출로 변환됨 ( 위 코드 참조 ) 

 

 

 

 

CUDA API 동기화 동작에는 2가지가 있다. 

 

1. Asynchoronous ( 비동기 )  : CPU 가 gpu 한테 명령 보내고 기다리지 않고 바로 다음줄 실행 ( <<< >>> 역시 여기 포함 ) 

2. synchoronous ( 동기 ) : GPU 가 끝날때까지 멈춰서 기다리는 것 

- cudaMemcpy ( D -> H ) 가 여기에 해당한다. 

 

(추가) 

Explicit synchronization : 비동기 작업 ㅎ웨 cpu 가 결과를 써야 하는 경우 cudaDeviceSynchronize()를 호출해서 수동으로 GPU 작업 완료를 기다리는 것. 

예를들어서 시간 측정하는 경우에 cpu 에서 기다렸다가 시간을 측정해야 하기 때문에 , cpu 를 명시적이 코드를 사용해서 멈춰두는것. 

 

 

( 따로 그림을 첨부하지는 않고 설명 ) 

 < GPU Initilization > 

 

왜 첫번째 커널 호출이 느린가 ? ( 과제할때 시간 측정을 해보면 실제로 첫 커널 실행은 다른 커널에 비해 매우 느리다. ) 

NVIDIA Nsight 프로파일링 결과들을 확인해보면 

 

- WARM_UP 의 경우 첫 커널이 cuLibraryLoadData 같은 초기화 때문에 오래 걸림 

- gpu 처음쓸 때 드라이버 초기화 , 컨텍슽트 생성 , 라이브러리 로딩 등 땜 느리다 ! 

이정도만 참고하면 된다. 

 

 

 

 

Warp Occupancy Revisited 

 

다시 warp occupacy 에 대해서 복습을 하고 이번 강을 마치도록 하자.

 

Warp occupancy 란 ? 

warp occupancy = Sm 에 현재 active 한 warp 수 / SM이 가질 수 있는 최대 warp 수 

 

TB가 SM에 배정되려면?

한 Thread Block(TB)이 SM에 올라가려면, 그 SM이 TB가 필요로 하는 레지스터와 shared memory를 제공할 수 있어야 한다. 자원이 허용하는 한도 내에서 여러 TB가 동시에 한 SM에서 실행될 수 있고. 

이전에 4가지 정도 제한을 잘 살펴봐야한다고 했는데 , 그런 이유들로 TB 이 SM 에 배정되는 경우가 다르다. 

따라서 엔지니어는 이를 잘 고려해서 배치해야한다. 

 

다만, OCCUPANCY 를 100%까지 맞출 필요는 없고 ( 불가능하고 ) , 적당히 약 70%정도 까지는 체워야 성능이 좀 나온다고 한다. 

 

 

예제를 몇가지 살펴보자. 

 

예제 1: 32x32 크기의 2D TB 2개를 한 SM에 동시에 실행 가능할까?

 

-> 32 X 32 (1024 ) thread 이게 2개면 , 2048 thread 인데 , cc 8.6 에서는 보면 SM 당 최대 1536 개 제한이 있으므로 불가능하다. 

 

예제 2: 256 크기의 1D TB은 최대 몇 개? 

256 / 32(warp 단위 ) = 8개warp 가 있다. ( cf. 하나의 thread block 은 하나의 SM 에 배정되어야 한다.  ) 

 

이제 세가지 조건을 살펴보자. 

- 블럭 수 제한 : 16개 -> 16개 까지 들어갈 수 있음. 

- WARP 수 : 48개 -> 6개 까지 들어갈 수 있다. 

- THREAD 수 제한 : 1536 / 256 = 6개 

 

위 3개의 조건에서 최소값이 6개 이므로 -

A : 6개의 thread block 이 들어갈 수 있다. 

그러면 총 48개의 warp 가 상주 할 수 있으므로 ( block 당 warp 수가 6개 )

-> 48/48 이다. ( 100 % ) 

 

나머지 예시 계산 확인용. 

1) TB size 16x16, regs/thread = 32, shared memory = 2048B

16x16 = 256 threads니까 block당 warp는 여전히 8개야.

• thread limit → 6 blocks
• warp limit → 6 blocks
• reg usage: 256 × 32 = 8192 regs/block → 65536 / 8192 = 8 blocks
• smem usage: 2048B/block → shared memory는 훨씬 여유 있음

결국 병목은 thread/warp 제한이므로 6 blocks, 48 warps, occupancy는 48/48 = 1.0. 슬라이드 결과와 맞아.

2) TB size 16x16, regs/thread = 40, shared memory = 2048B

• block당 threads = 256
• block당 warps = 8
• reg usage: 256 × 40 = 10240 regs/block
• 65536 / 10240 = 6 blocks

( 여기는 오타가 좀 있는듯 ) -> occpancy 1 이 맞는거 것 같음.  

 

3) TB size 16x16, regs/thread = 41, shared memory = 2048B

여기가 진짜 재밌는 부분이야. 41개로 딱 1개 늘었는데 block 수가 6개에서 5개로 떨어진다. 이유는 register allocation granularity 때문이야. 슬라이드 계산에도

CEILING(41*32, 256) = 1536
...
결과적으로 register 제한으로 5 blocks

 

라고 나와 있어. 즉 warp 단위/정렬 단위 때문에 register 사용량이 계단식으로 튀어 올라간다. 그래서 41 regs/thread에서는 5 blocks, 40 warps가 resident하고, occupancy는 40/48 = 0.833이 된다. 이 값은 슬라이드의 TB 5 / W 40 (Occupancy = 0.833)와 정확히 맞아

 

 

--> 256단위로 배정하는건 어디 ? 

 

Register allocation unit size -  256

Register allocation granularity - warp 

 

 

• 레지스터는 warp 단위로 할당된다 (thread 단위가 아님)
• 할당할 때 256개 단위로 올림해서 배정한다 
  
  
  

  ---

  ① 블록 크기 제한② 레지스터 제한 ← 여기가 핵심41×32 = 1312인데, 256 단위로 올림하니까 1536이 됨. reg=40일 때는 1280이었는데, 1개 늘었을 뿐인데 256 올라감.4 단위로 내림하니까 42→40.③ Shared Memory 제한④ 최종 결과
  제한 요소최대 TB 수

  블록/워프 크기	6
  레지스터	5 ← 병목
  Shared Memory	32

   

   

  TB = 5개
  Warps = 5 × 8 = 40
  Occupancy = 40 / 48 = 0.833 (83%)

  reg=40 vs reg=41 비교레지스터 1개 차이인데, 256 단위 할당(allocation unit size) 때문에 warp당 256개가 추가로 잡히면서 TB가 하나 빠지고, occupancy가 17%나 떨어지는 거야. 이게 계단식 변화의 대표적인 예시.
• reg=40reg=41

  warp당 할당 레지스터	1280	1536
  SM 최대 warp	48	40
  최대 TB	6	5
  Occupancy	1.0	0.833

• 세 제한 중 가장 작은 값이 적용됨:
•  

   

  MySharedMemPerBlock = CEILING(2048, 128) = 2048
  최대 TB 수 = FLOOR(65536/2048, 1) = 32  ← 여유로움

•  

   

  레지스터 기준 최대 TB 수 = FLOOR(40/8, 1) × FLOOR(65536/65536, 1)
                           = 5 × 1 
                           = 5

•  

   

  SM 전체 최대 warp 수 = FLOOR(65536 / 1536, 4) 
                        = FLOOR(42.67, 4) 
                        = 40

•  

   

  warp당 레지스터 할당량 = CEILING(41 × 32, 256) 
                         = CEILING(1312, 256) 
                         = 1536

•  

   

  warps/block = CEILING(256/32, 1) = 8
  블록 수 제한 = MIN(16, FLOOR(48/8, 1)) = MIN(16, 6) = 6

• 기본 조건: TB size 16×16 = 256 threads, 8 warps/TB, smem 2048 bytes, CC 8.6

 

 

 

 

 

역시 그저 계산과 공식인 영역이라. 나중에 필요할때 참고하면 된다. 

중요한건 warp 을 얼마나 잘 배정할것이냐. 이걸 어떻게 분배해야 적당한 값으로 최대의 효율을 낼 수 있는가의 문제이다.