kernel_hierarchy_test

grid - block - warp - thread(bank)  계층 이해하기 위한 test 코드

        const int rows = 8;
        const int cols = 8;

        dim3 blockDim(8, 8); // 블록당 64 threads → warp 2개
        dim3 gridDim((cols + blockDim.x - 1) / blockDim.x, (rows + blockDim.y - 1) / blockDim.y);

        int num_elems = rows * cols;

        ThreadInfo* d_info = nullptr;
        ThreadInfo* h_info = new ThreadInfo[num_elems];

        cudaMalloc(&d_info, num_elems * sizeof(ThreadInfo));

        hierarchy_test_kernel<<<gridDim, blockDim>>>(d_info, rows, cols);
        cudaDeviceSynchronize();

        cudaMemcpy(h_info, d_info,
                   num_elems * sizeof(ThreadInfo),
                   cudaMemcpyDeviceToHost);

 

const int cols, rows = 8; 호스트 변수, 문제 크기 설정

 

dim3 blockDim(8, 8);

• CUDA 에서 블록의 스레드 배치르르 나타내는 타입, dim3
• blockDim.x = 8, blockDim.y = 8, blockDim.z = 1 
• 블록 당 스레드 수 = 64

이 블록의 64개 스레드는 커널 안에서 쓰는 threadIdx.x, threadIdx.y 값이 0~7 범위를 돈다는 뜻

warp 관점에서 보면, warpSize = 32, 블록당 2 warps 

 

dim3 gridDim((cols + blockDim.x - 1) / blockDim.x, (rows + blockDim.y - 1) / blockDim.y );

이건 전체 문제를 몇 개의 블록으로 쪼갤지를 정하는 부분

각 축별로, 

x 방향 블록 개수 (8 + 8 - 1) / 8 = 1

y 방향도 동일, 최종적으로 블록이 1개

blockDim.x - 1 패턴은 보통 ceil 나눗셈을 하기 위해 쓰는 템플릿.

• grid 에는 1개의 block
• block 안에 thread 64개

 

int num_elems = rows * cols;

이번 테스트에서 다루는 전체 셀의 개수

각 global_row, global_col 쌍마다 ThreadInfo 하나씩 저장, 

 

ThreadInfo* d_info = nullptr;

ThreadInfo* h_info = new ThreadInfo[num_elems];

d_info : 디바이스 메모리 포인터용 변수, 

h_info : 호스트 메모리에 할당, 길이 num_elems = 64 인 ThreadInfo 배열, 나중에 cudaMemcpy 로 GPU 에서 결과를 가져온 후 여기에 채워서 로그 출력

 

cudaMalloc(&d_info, num_elems * sizeof(ThreadInfo));

여기서 비로소 GPU 메모리 할당이 일어난다. 

두 번째 인자 : 필요한 바이트 수, 64 * sizeof(ThreadInfo)

cudaMalloc 의 첫 번째 인자는 포인터 변수의 주소, 내부에서 GPU 메모리 주소를 만들어서 d_info 에 집어넣음

 

d_info 는 GPU 글로벌 메모리 상의 배열시작 주소를 가리키게 됨, 

ThreadInfo d_info[64]; 를 글로벌 메모리에 만들어준 느낌

 

 

hierarchy_test_kernel<<<gridDim, blockDim>>>(d_info, rows, cols); 

여기가 진짜 GPU 실행 요청

<<<gridDim, blockDim>>>

드라이버/ 런타임에게

gridDim.x = 1, gridDim.y = 1 -> 블록 1개

blockDim.x =8, blockDim.y = 8 -> 블록당 스레드 64개

위 구성을 가진 커널을 실행하라고 요청

 

논리적으로 생성되는 스레드

(blockIdx.x, blockIdx.y) = (0, 0) 블록 한 개

그 안에 (threadIdx.x, threadIdx.y) = (0~7, 0~7) 스레드들, 64개

 

각 스레드는 코드 한 번씩 독립적으로 실행

__global__ void hierarchy_test_kernel(ThreadInfo* info,
                                      int rows, int cols)
{
    ...
}

이 함수의 본문이 스레드 하나 기준 실행 흐름

 

d_info, rows, cols 는 각각 

d_info : GPU 글로벌 메모리 배열 포인터

rows, cols : 8 으로 각 스레드에게 전달됨

 

이 함수 호출은 CPU 입장에서는 비동기 asynchronous, 

 

cudaDeviceSynchronize();

이 호출은 CPU 에서 GPU 를 기다리는 함수 

지금까지 큐에 들어간 모든 커널/작업이 다 끝날 때까지 여기서 블록하고 대기

hierarchy_test_kernel<<<...>>> 을 날렸음, 그 커널이 다 끝날 때까지 기다렸다가 GPU 의 끝난 신호를 보내면 함수 리턴

 

cudaMemcpy(... cudaMemcpyDeviceToHost) 자체도 동기 함수라 내부적으로  GPU 작업 완료를 기다린다.

 

 

cudaMemcpy(h_info, d_info, num_elems * sizeof(ThreadInfo), cudaMemcpyDeviceToHost);

이제 GPU 에서 CPU 로 결과를 가져오는 단계.

from d_info, to h_info

바이트 수 : 64 * sizeof(ThreadInfo)

방향 플래그 : cudaMemcpyDeviceToHost

GPU 에서 해당 메모리 범위 읽기 완료

CPU 메모리에 복사 완료

h_info[i]를 자유롭게 사용 가능

 

다음으로 hierarchy_test_kernel 확인

함수 시그니처

__global__ void hierarchy_test_kernel(ThreadInfo* info, int rows, int cols)

• __global__ : 이 함수는 GPU 에서 실행되고, CPU 에서 <<<...>>> 로 호출되는 커널 함수라는 뜻
• ThreadInfo* info : GPU 글로벌 메모리에 있는 ThreadInfo 배열 시작 주소
• int rows, int cols : 호스트에서 넘겨준 정수 값

스레드 관점에서 info 라는 배열 포인터, rows, cols = 8 을 전달받고 실행 시작

 

2D 전역 인덱스 계산

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

전체 문제를 rows * cols 행렬이라고 생각하고 각 스레드에게 그 중 한 칸씩 맡으려는 상황

블록/스레드 2D 좌표 -> 전체 좌표로 바꾸는 공식

global_row = 0 * 8 + 2 = 2;
global_col = 0 * 8 + 3 = 3;

이 스레드는 전역 좌표 (2,3) 셀을 맡는다고 해석

출력에서의 모습

[global ( 2, 3)]  block(0,0)  thread(3,2)  ...

이 라인이 바로 이 계산 결과를 찍어준 것

전체 8*8 문제 중에서 몇 번째 row/col 을 맡았는가

 

범위 체크를 통한 문제 크기를 넘어가는 스레드 제외

if (global_row >= rows || global_col >= cols) return;

 

 

블록 내 선형 thread 인덱스

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

2D threadidx -> 1D 인덱스 변환

• blockDim.x = 8
• row 마다 8개씩 스레드 존재
• tid_in_block = y * 8 + x 로 flatten

 

warp / lane 계산

int warp_size = warpSize; // 일반적으로 32
int warp_in_block = tid_in_block / warp_size;
int lane_in_warp = tid_in_block % warp_size;

warp_size

• 거의 항상 32
• 하드웨어에서 한 번에 같이 실행하는 스레드 묶음의 크기

warp_in_block

warp_in_block = tid_in_block / 32;

• 블록 안에서 몇 번째 warp 에 속하는지.
• 0 번 warp : tid 0~31
• 1 번 warp : tid 32~63

 

 

전역 선형 인덱스 (info 배열 인덱스)

int idx = global_row * cols + global_col;

• 이제는 블록안이 아닌 전체 문제 기준 flatten
• 전체를 rows * cols 2D 그리드라고 보면,
  • 한 row 에 cols 개씩 존재

 

ThreadInfo 구조체 채우기

ThreadInfo ti;
ti.block_x       = blockIdx.x;
ti.block_y       = blockIdx.y;
ti.thread_x      = threadIdx.x;
ti.thread_y      = threadIdx.y;
ti.warp_in_block = warp_in_block;
ti.lane_in_warp  = lane_in_warp;
ti.global_row    = global_row;
ti.global_col    = global_col;

 

 

글로벌 메모리에 쓰기

info[idx] = ti;

• info 는GPU 글로벌 메모리에 있는 ThreadInfo 배열 주소
• idx 는 아까 계산한 전역 1D 인덱스

전역 좌표를담당하는 스레드가 info[global_row * cols + global_col] 위치에 자기 정보를 써 넣는다.