Configurable battalion — block × thread 손가락 끝으로 느끼기
이전 <<<1, 4>>> launch는 4마리 분대 한 개. 현실 launch는 수백 block × block당 수백 thread. 아래 확장판은 command line에서 block 수랑 thread 수 받아서 어떤 대형이든 시도하고 무슨 일 일어나는지 볼 수 있게 해줘:
~12 min · cuda, hello, battalion, global-id
이전 <<<1, 4>>> launch는 4마리 분대 한 개. 현실 launch는 수백 block × block당 수백 thread. 아래 확장판은 command line에서 block 수랑 thread 수 받아서 어떤 대형이든 시도하고 무슨 일 일어나는지 볼 수 있게 해줘:
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
__global__ void ant_battalion_report() {
int ant_id = blockIdx.x * blockDim.x + threadIdx.x;
printf("Ant %d 보고, block %d, slot %d.\n",
ant_id, blockIdx.x, threadIdx.x);
}
int main(int argc, char **argv) {
if (argc != 3) {
fprintf(stderr, "Usage: %s <num_blocks> <threads_per_block>\n", argv[0]);
return 1;
}
int blocks = atoi(argv[1]);
int threads = atoi(argv[2]);
if (threads % 32 != 0) {
fprintf(stderr, "Tip: threads-per-block은 32 (warp size)의 배수가 좋아.\n");
}
ant_battalion_report<<<blocks, threads>>>();
cudaDeviceSynchronize();
return 0;
}nvcc -arch=sm_89 ant_battalion.cu -o ant_battalion
./ant_battalion 1 4 # 4마리 분대 1개 — 원래 hello랑 동일
./ant_battalion 2 8 # 분대 2 × 8 = 16마리 — 순서 interleave
./ant_battalion 32 32 # 1024마리 — scheduler 혼돈 첫 힌트
./ant_battalion 128 4 # 작은 분대 많이로 512마리
# 주목:
# 1) 개미가 숫자 순으로 print 안 함 — 스케줄러가 block 자유롭게 픽.
# 2) block 안에선 threadIdx 값들이 보통 그룹화됨.
# 3) ~10K thread 넘으면 printf throughput 급감 — buffer 참.ant_battalion.cu 컴파일하고 ./ant_battalion 4 32, ./ant_battalion 32 32, ./ant_battalion 256 256 셋 다 돌려. 각각 line 수 세서 (block × thread랑 같아야 함) 높은 count에서 'random' 하게 보이는 정도 인지. 그 다음 파일로 redirect: ./ant_battalion 256 256 > out.txt, 특정 ant_id 하나 grep — 출력 순서 혼돈해도 GPU가 정확한 거 확인.아직 댓글이 없어요. 첫 댓글을 남겨보세요.