제2부 CUDA 프로그래밍

2.3 실행 구성 — 그리드·블록·grid-stride 설계

CUDA 커널을 런치할 때 <<<grid, block>>> 안에 무엇을 넣을지가 이 장의 전부다. 숫자 두 개처럼 보이지만, 이 선택이 점유율·메모리 병합·다이버전스·런치 오버헤드를 한꺼번에 결정한다. "32의 배수로 256 안팎"은 절반만 맞다 — 그 절반을 끝까지 채운다.

2.3.1 이 장에서 배우는 것 / 선수 지식

배우는 것:

• 1D·2D 데이터에 스레드를 매핑하고 전역 인덱스를 계산하는 관용구를 쓴다.
• 블록 크기를 워프(32) 배수로 설정하는 이유를 점유율·다이버전스·자원 한계로 설명한다.
• 블록 크기 선택(128·256)의 근거와, "점유율 최대 ≠ 성능 최대"를 안다.
• 워프 다이버전스를 일으키는 코드 패턴을 알아보고 회피한다.
• grid-stride loop로 그리드 크기를 데이터에서 분리하고, 점유율 API로 출발 구성을 자동 산출한다.

선수 지식: 2.1~2.2(런치 문법·인덱싱·grid-stride·검증), 1.2(SIMT·워프), 1.4~1.5(SM 내부·스케줄링). 점유율의 정의·수계산·점유율 vs ILP는 1.6에서 다뤘다 — 이 장은 그것을 블록 크기 선택에 적용한다(반복하지 않고 참조).

2.3.2 실행 구성은 왜 설계 대상인가

커널 런치 kernel<<<grid, block, shmem, stream>>>(...)의 처음 두 인자가 실행 구성이다. block(=blockDim)은 한 블록의 스레드 수·모양, grid(=gridDim)은 블록의 수·모양(둘 다 dim3, 최대 3차원). 이 두 숫자를 고르는 것이 "설계"인 이유는 하드웨어가 이 논리 구성을 물리 자원에 매핑하는 방식 때문이다. 세 사실이 모든 설계를 지배한다:

1. 블록 전체가 단일 SM에 배정되고, 끝날 때까지 그 SM을 떠나지 않는다(여러 SM에 분할되지 않음). 따라서 블록 하나가 쓰는 자원(레지스터·공유 메모리)이 SM에 몇 개 들어가는지를 결정한다(1.6).
2. 실행 단위는 블록이 아니라 워프(32 스레드)다. 블록은 32개 단위로 잘려 워프가 된다. 이것이 블록 크기를 32의 배수로 설정하는 근본 이유.
3. SM은 여러 워프를 동시에 상주시켜 두고 교체해 지연을 숨긴다(1.2·1.5). 상주 워프가 많을수록 숨길 지연이 많다 — 단 항상 좋은 것은 아니다(1.6).

2.3.3 인덱싱 — 스레드를 데이터에 매핑

2.3.3.1 1D 매핑

전역 인덱스 = blockIdx.x * blockDim.x + threadIdx.x. 런치 시 블록 수는 올림 나눗셈 (n+block-1)/block(2.2).

2.3.3.2 2D 매핑 (이미지·행렬)

__global__ void invert(unsigned char* img, int w, int h) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (col < w && row < h) {
        int idx = row * w + col;          // 행 우선(row-major) 선형화
        img[idx] = 255 - img[idx];
    }
}
dim3 block(16, 16);                       // 256 스레드 = 8 워프
dim3 grid((w + 15) / 16, (h + 15) / 16);
invert<<<grid, block>>>(img, w, h);

왜 threadIdx.x가 열(col)이어야 하는가: 다차원 블록도 하드웨어엔 선형화된 1D로 보인다 — linear = threadIdx.x + threadIdx.y*blockDim.x + .... 이 linear을 32개씩 끊은 것이 워프이므로 threadIdx.x가 워프를 채우는 빠른 축이다. 행 우선 배열에서 같은 행의 인접 열은 메모리에서 인접하므로, threadIdx.x→col로 매핑해야 워프 내 32 스레드가 연속 주소를 읽어 병합 접근(coalescing)이 된다(1.8). x를 row로 잘못 매핑하면 워프가 w 간격으로 흩어진 주소를 읽어 대역폭이 무너진다.

linear = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y
warp_id = linear/32,  lane = linear%32
→ threadIdx.x 가 빠른 축. 2D에선 x를 "메모리 연속 축(보통 열)"에.

2.3.4 블록 크기 선택 — 32 정렬과 워프 배수

블록 크기는 항상 warpSize(=32)의 배수로. 이유 셋:

1. 낭비 없는 워프: 32의 배수가 아니면 마지막 워프가 부분적으로 채워진다(블록 48 → 워프 2개 점유, 16 레인 비활성). 그 분의 실행 자원·스케줄 슬롯이 낭비.
2. 점유율 정렬: SM의 워프 슬롯은 정수 단위. 부분 워프면 SM 한도를 어중간하게 채워 점유율이 떨어진다(1.6).
3. 다이버전스·병합 정렬: 워프 경계가 데이터 경계와 어긋나면 분기·메모리 접근이 워프를 가로질러 깨진다.

128·256 중 무엇?

• 하한: 너무 작으면(32·64) SM의 블록 수 한도(보통 SM당 16~32 블록, 확인 필요)에 먼저 걸려 워프를 다 못 채운다. 블록당 고정 오버헤드도 상대적으로 커진다.
• 상한: 너무 크면(1024) 블록 하나가 SM 자원을 독점해 SM당 블록이 1~2개로 줄고 레지스터·공유 메모리 한계에 일찍 부딪힌다. __syncthreads가 1024 스레드를 모두 기다려 지연도 커진다.
• 스윗스팟: 128 또는 256이 대부분 좋은 출발점. 워프 4~8개로 지연 은닉에 충분하고 SM당 여러 블록이 들어가 스케줄 유연성이 생긴다.

견해: "모르겠으면 256으로 시작하고, 레지스터/공유 메모리 많이 쓰는 커널은 128로 낮춰 더 많은 블록이 상주하게 한다. 그다음 프로파일러(2.14)로 확정한다." 출발점일 뿐 최적값은 커널·아키텍처마다 다르다. 블록당 스레드는 하드웨어 한도(보통 1024, 확인 필요) 를 넘을 수 없다.

블록 크기가 SM당 동시 블록·워프·점유율을 어떻게 정하는지(레지스터·공유 메모리·하드 한도의 최솟값)는 1.6에서 수계산했다. 블록 크기를 바꿔가며 그 계산을 적용하면 "왜 이 크기?"의 정량적 근거가 나온다.

2.3.5 점유율과 성능 — 더 높은 점유율이 늘 빠른가

아니다(1.6에서 정밀하게 다룸). 점유율은 지연을 숨길 워프가 충분한가의 척도일 뿐, 성능과 단조 증가 관계가 아니다.

• 낮은 점유율: 메모리 대기를 숨길 워프가 없어 SM이 유휴 상태다 → 상주 워프를 늘리면 빨라진다.
• 충분 지점 이후: 지연이 이미 숨겨졌다 → 워프를 더 상주시켜도 거의 안 변한다(보통 50~70%면 충분).
• 때로 낮은 점유율이 더 빠르다: 레지스터를 넉넉히 써서 데이터를 레지스터에 들고 있으면 메모리 왕복이 줄어 ILP로 지연을 숨긴다(1.6의 점유율 vs ILP).

견해: 점유율은 진단 지표이지 최적화 목표가 아니다. "100% 만들기"를 좇아 레지스터를 줄이면 스필로 더 느려질 수 있다. 순서: ① 짠다 → ② 프로파일러(2.14)로 실제 병목을 본다 → ③ 점유율이 병목일 때만 블록 크기·레지스터를 손본다. 메모리 바운드 판정은 1.10.

2.3.6 워프 다이버전스 회피

한 워프의 32 스레드는 같은 명령을 락스텝 실행한다. if로 스레드마다 다른 경로로 갈리면 하드웨어는 두 경로를 순차 실행하고 해당 안 되는 레인을 마스킹한다(다이버전스, 1.2). 분기 자체가 문제가 아니라, 분기가 워프 경계를 가로지를 때만 문제다.

회피·완화:

• 워프 단위로 갈리게 데이터 정렬·배치: 조건이 참인 원소를 모아두면 워프 전체가 한 방향.
• 워프 정렬 인덱스로 분기: if ((i/32)%2)는 워프 단위라 다이버전스 없음. if (i%2)는 레인 단위라 발생.
• 분기를 산술로 대체(predication): 짧은 if는 컴파일러가 조건부 실행으로 바꿔 직렬화를 피한다. y = cond ? a : b;. 단 양쪽 다 계산하므로 본문이 부하가 크면 손해.
• 경계 검사(if (i<n))는 보통 무해: 마지막 워프 하나만 갈리고 나머지는 전부 참. 안전이 우선이니 빼라는 뜻이 아니다.

2.3.7 grid-stride loop와 점유율 API

grid-stride loop(2.2)는 그리드 크기를 N에서 분리한다. 그리드 크기는 점유율 API로 정하면 간결하다:

int blockSize = 0, minGridSize = 0;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, saxpy, 0, 0);
//   인자: 동적 공유메모리 바이트(0), blockSizeLimit(0=제한없음)
saxpy<<<minGridSize, blockSize>>>(n, 2.0f, x, y);   // grid-stride라 minGridSize 그대로

// 비-stride 커널이면 N으로 올림:
int gridForN = (n + blockSize - 1) / blockSize;

관련: cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, kernel, blockSize, dynamicSMem) — 특정 블록 크기에서 SM당 동시 블록 수(1.6 수계산의 런타임판). 이론 점유율 = numBlocks*blockSize/maxThreadsPerSM.

비판적으로: 이 API는 이론 점유율을 최대화하는 블록 크기를 준다. 점유율 최대 ≠ 성능 최대(2.3.5). 따라서 출력은 좋은 출발점이지 최종값이 아니다. 메모리 바운드 커널은 추천값이 잘 맞지만, 레지스터를 의도적으로 많이 쓰는(낮은 점유율·높은 ILP) 커널은 직접 튜닝이 낫다. 항상 프로파일러로 확정.

2.3.8 중간 점검 (체크포인트)

• [ ] 1D·2D 전역 인덱스를 외워 쓰고, 2D에서 threadIdx.x를 메모리 연속 축(열)에 두는 이유를 병합으로 설명한다.
• [ ] 블록 크기를 항상 32의 배수로, 출발점을 128·256으로 설정하는 이유를 제시한다.
• [ ] "점유율 100%가 항상 가장 빠른가?"에 아니다라고 답하고 충분 지점·ILP 사례를 든다(1.6).
• [ ] if (i%2)와 if ((i>>5)&1) 중 어느 것이 다이버전스를 일으키는지 즉답한다(전자).
• [ ] grid-stride의 stride(gridDim.x*blockDim.x)와 점유율 API의 한계를 제시한다.

확인 신호: 블록 256·128 성능이 거의 같다면 → 이미 충분 점유율이거나 메모리 바운드(1.10). 블록을 48·100으로 작성했다면 → 워프가 부분 채워져 점유율이 샌다, 32 배수로 교정.

2.3.9 흔한 함정

• 32 비배수 블록(48·100 등). 마지막 워프가 부분 채워져 자원이 샌다. 반드시 32 배수.
• 점유율만 좇기. 100%를 목표로 레지스터를 줄이다 스필 → local memory(글로벌만큼 느림) → 더 느려짐(1.6).
• 2D 매핑에서 x·y 축 혼동. threadIdx.y를 열에 매핑하면 워프가 띄엄띄엄 읽어 병합이 깨지고 대역폭 수배 손실. x=연속 축.
• 블록 1024 근거 없는 사용. 가장 큰 것이 가장 좋다는 착각. SM당 블록이 1~2개로 줄어 스케줄 유연성·__syncthreads 효율이 나빠질 수 있다.
• 레인 단위 분기 남발. i%2, i%32<k 같은 워프 내부 조건이 다이버전스. 워프 단위 정렬·predication으로.
• 그리드 크기를 N에 강결합. 거대 N에 1:1 매핑으로 그리드 폭발. grid-stride로 분리.
• 점유율 API 출력 맹신. 이론 점유율 최대 ≠ 성능 최대. 출발점으로만.

2.3.10 연습문제

연습 1 (1D 런치 구성). 길이 n=1,000,003(소수) 배열에 SAXPY를 블록 256 비-stride로. 그리드 크기와 남는 스레드 수는? 경계 가드는 왜 필요한가?

연습 2 (다이버전스 판별). 다음 중 다이버전스를 일으키는 것을 고르고 이유를(블록 256): (a) if (blockIdx.x%2) (b) if (threadIdx.x%2) (c) if ((threadIdx.x/32)%2).

연습 3 (2D 매핑·병합). 1920×1080 이미지에 16×16 블록으로 런치하는 그리드는? threadIdx.x를 row에 매핑하면 row-major 저장 기준 무엇이 잘못되나?

연습 4 (낮은 점유율 고성능). 어떤 행렬곱 커널이 점유율 50%인데 점유율 100% 구현보다 빠르다. 어떻게 가능한가? 무엇을 측정해 확인하나?

2.3.11 요약 · 다음 장

• 실행 구성은 숫자 두 개가 아니라 설계다: 블록 전체가 단일 SM에 배정, 실행 단위는 워프(32), SM은 상주 워프로 지연을 숨긴다.
• 블록 크기는 항상 32의 배수, 출발점은 128·256. SM당 동시 블록·점유율 수계산은 1.6.
• 점유율은 진단 지표이지 목표가 아니다. 다이버전스는 분기가 워프 경계를 가로지를 때만 문제 — 워프 단위 정렬·predication으로 회피.
• grid-stride로 그리드 크기를 N에서 분리하고, 점유율 API로 출발 구성을 산출하되 출력은 출발점으로만.

다음 장: 2.4 메모리 모델 I — 종류. 글로벌·상수·텍스처·로컬의 특성과 용도를 본다.