CUDA Vector Addition - GPU 병렬 처리 기초

__global__ void addKernel(int* c, const int* a, const int* b) { int i = threadIdx.x; // "나는 몇 번째 스레드인가?" c[i] = a[i] + b[i]; }

__global__ void addKernel(...)

__global__ // GPU에서 실행, CPU에서 호출 __device__ // GPU에서 실행, GPU에서 호출 __host__ // CPU에서 실행, CPU에서 호출 (기본값)

for (int i = 0; i < 5; i++) { c[i] = a[i] + b[i]; } // 총 5번의 반복, 순차 실행

// 5개의 스레드가 동시에 실행 // Thread 0: c[0] = a[0] + b[0] // Thread 1: c[1] = a[1] + b[1] // Thread 2: c[2] = a[2] + b[2] // Thread 3: c[3] = a[3] + b[3] // Thread 4: c[4] = a[4] + b[4]

int i = threadIdx.x;

threadIdx.x // 블록 내 x축 스레드 인덱스 (0부터 시작) threadIdx.y // 블록 내 y축 스레드 인덱스 threadIdx.z // 블록 내 z축 스레드 인덱스 blockIdx.x // 그리드 내 x축 블록 인덱스 blockIdx.y // 그리드 내 y축 블록 인덱스 blockIdx.z // 그리드 내 z축 블록 인덱스 blockDim.x // 블록의 x축 크기 (스레드 개수) gridDim.x // 그리드의 x축 크기 (블록 개수)

// 1D 예시 int i = blockIdx.x * blockDim.x + threadIdx.x; // 2D 예시 (이미지 처리) int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int idx = y * width + x;

cudaMalloc((void**)&dev_a, size * sizeof(int));

cudaError_t cudaMalloc(void** devPtr, size_t size); // devPtr: 할당된 GPU 메모리 주소를 받을 포인터 // size: 할당할 바이트 크기 // 반환값: 성공 시 cudaSuccess

int* dev_a; // GPU 메모리 주소를 저장할 포인터 cudaMalloc((void**)&dev_a, 100 * sizeof(int)); // 100개 int 공간

cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);

cudaError_t cudaMemcpy( void* dst, // 목적지 주소 const void* src, // 소스 주소 size_t count, // 복사할 바이트 크기 enum cudaMemcpyKind kind // 복사 방향 );

cudaMemcpyHostToHost // CPU → CPU cudaMemcpyHostToDevice // CPU → GPU cudaMemcpyDeviceToHost // GPU → CPU cudaMemcpyDeviceToDevice // GPU → GPU cudaMemcpyDefault // 자동 감지 (Unified Memory)

addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

kernelFunction<<<gridDim, blockDim, sharedMem, stream>>>(args...); // gridDim: 그리드 크기 (블록 개수) // blockDim: 블록 크기 (블록당 스레드 개수) // sharedMem: 공유 메모리 크기 (바이트, 선택) // stream: CUDA 스트림 (선택)

addKernel<<<1, size>>>(dev_c, dev_a, dev_b); // <<<1, size>>>의 의미: // - 1개의 블록 생성 // - 각 블록에 size(5개)만큼의 스레드 생성 // - 총 1 × 5 = 5개의 스레드

// 1D 구성 kernel<<<numBlocks, threadsPerBlock>>>(); // 2D 구성 (이미지 처리) dim3 threadsPerBlock(16, 16); dim3 numBlocks(width/16, height/16); kernel<<<numBlocks, threadsPerBlock>>>(); // 3D 구성 (볼륨 데이터) dim3 threadsPerBlock(8, 8, 8); dim3 numBlocks(width/8, height/8, depth/8); kernel<<<numBlocks, threadsPerBlock>>>();

cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c);

for (int i = 0; i < N; i++) { c[i] = a[i] + b[i]; } // 시간 복잡도: O(N) // 1백만 개: ~1초

addKernel<<<blocks, threads>>>(dev_c, dev_a, dev_b); // 시간 복잡도: O(1) (이론적) // 1백만 개: ~0.001초

for (int y = 0; y < height; y++) { for (int x = 0; x < width; x++) { // 각 픽셀에 대해 광선 추적 Ray r = camera.getRay(x, y); Color c = trace(r, scene); image[y][x] = c; } } // 1920×1080 = 2,073,600 픽셀을 순차 처리 // 매우 느림 (수 분 ~ 수 시간)

__global__ void render(Color* image, Camera camera, Scene scene, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= width || y >= height) return; int idx = y * width + x; Ray r = camera.getRay(x, y); image[idx] = trace(r, scene); } // 각 픽셀마다 스레드 하나 배정 // 수천, 수만 개의 스레드가 동시에 계산 // 매우 빠름 (수 초 ~ 수 분)

// Vector Addition threadIdx.x → 배열 인덱스 i c[i] = a[i] + b[i] // Ray Tracing threadIdx.x, threadIdx.y → 픽셀 좌표 (x, y) image[y][x] = trace(ray(x, y))

cudaError_t err; err = cudaMalloc((void**)&dev_a, size * sizeof(int)); if (err != cudaSuccess) { fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(err)); return 1; }

#define CHECK_CUDA(call) \ do { \ cudaError_t err = call; \ if (err != cudaSuccess) { \ fprintf(stderr, "CUDA Error: %s at %s:%d\n", \ cudaGetErrorString(err), __FILE__, __LINE__); \ exit(1); \ } \ } while(0) // 사용 CHECK_CUDA(cudaMalloc((void**)&dev_a, size * sizeof(int)));

addKernel<<<1, size>>>(dev_c, dev_a, dev_b); // 커널 실행 에러 확인 cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { fprintf(stderr, "Kernel launch failed: %s\n", cudaGetErrorString(err)); } // 커널 완료 대기 및 런타임 에러 확인 err = cudaDeviceSynchronize(); if (err != cudaSuccess) { fprintf(stderr, "Kernel execution failed: %s\n", cudaGetErrorString(err)); }

__global__ void addKernel(int* c, const int* a, const int* b) { int i = threadIdx.x; printf("Thread %d: %d + %d = %d\n", i, a[i], b[i], a[i] + b[i]); c[i] = a[i] + b[i]; }

int device; cudaGetDevice(&device); cudaDeviceProp prop; cudaGetDeviceProperties(&prop, device); printf("Device: %s\n", prop.name); printf("Compute Capability: %d.%d\n", prop.major, prop.minor); printf("Max Threads per Block: %d\n", prop.maxThreadsPerBlock); printf("Max Grid Size: %d × %d × %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);

for (int i = 0; i < iterations; i++) { cudaMemcpy(dev_data, host_data, size, cudaMemcpyHostToDevice); kernel<<<blocks, threads>>>(dev_data); cudaMemcpy(host_data, dev_data, size, cudaMemcpyDeviceToHost); }

cudaMemcpy(dev_data, host_data, size, cudaMemcpyHostToDevice); for (int i = 0; i < iterations; i++) { kernel<<<blocks, threads>>>(dev_data); } cudaMemcpy(host_data, dev_data, size, cudaMemcpyDeviceToHost);

// 블록당 스레드 수는 32의 배수 (Warp 크기) // 일반적으로 128, 256, 512 사용 dim3 threadsPerBlock(256); dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x); kernel<<<numBlocks, threadsPerBlock>>>(data, N);

__global__ void slowKernel(int* data) { int idx = threadIdx.x; // Global Memory 반복 접근 (느림) data[idx] = data[idx] * 2 + data[idx] * 3; } __global__ void fastKernel(int* data) { __shared__ int temp[256]; int idx = threadIdx.x; // Global Memory에서 Shared Memory로 복사 temp[idx] = data[idx]; __syncthreads(); // Shared Memory 접근 (빠름) temp[idx] = temp[idx] * 2 + temp[idx] * 3; __syncthreads(); // 결과를 Global Memory에 저장 data[idx] = temp[idx]; }