[CUDA] 벡터 합 예제

해당 포스트에서는 CUDA를 이용하여 어떻게 벡터합을 계산하는지를 고려한다.

두 벡터의 합

1. Process

(1) cudaMallc()을 호출하여 입력 배열 dev_a, dev_b, dev_c에 대한 공간을 각각 디바이스에 할당한다. 
(2) cudaMemcpy()를 사용하여
    - cudaMemcpyHostToDevice: 입력데이터를 디바이스로 복사 
    - cudaMemcpyDeviceToHost: 결과데이터를 호스트로 복사 
(3) 모두 사용한 Device의 메모리를 해제한다. 

2. code

#include <stdio.h>
#include "cuda_runtime.h"

#define N 10


__global__ void add(int *a, int *b, int *c){
	int tid = blockIdx.x;
    // Device에 N개의 블록을 실행시켰으므로, 
    // block의 index가 N을 넘지 않도록 체크를 해야한다. (Block index의 범위: 0~N-1)
    // 만일, 해당 부분을 체크하지 않으면, 할당되지 않은 메모리의 Dummy values들이 사용될 
    // 수 있다. 
	if (tid < N){
		c[tid] = a[tid] + b[tid];
        
	}
};
int main(void){
	
	int a[N], b[N], c[N];
	int *dev_a, *dev_b, *dev_c;

	// allocation of device memory
	cudaMalloc( (void **)&dev_a, N * sizeof(int));
	cudaMalloc( (void **)&dev_b, N * sizeof(int));
	cudaMalloc( (void **)&dev_c, N * sizeof(int));

	for (int i = 0; i <N; i++){
		a[i] = -i;
		b[i] = i * i;
	}

	// copy array of a and b into device memory 
	cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy( dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);

	add<<<N,1>>>(dev_a, dev_b, dev_c);

	// copy array of dev_c into Host memory

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

	for (int i = 0; i < N; i ++){
		printf( "%d + %d = %d \n", a[i], b[i], c[i]);
	}

	// free memory of device

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

 

※ Note
(1) 용어 정리 
    - 블록 (block): 병렬 실행된 각각의 코드 
    - 그리드 (grid): 병렬 블록의 집합 
(2) __global__ void add() 함수 내의 변수 blockIdx.x 
   - 위 변수 blockIdx는 cuda runtime에서 이미 내장 변수(built-in)로 정의가 되어있기 때문에, 별도로 정의를 할 필요가 없다. 
   - 위 변수를 이용하여 현재 디바이스 코드를 작동하는 블록의 인덱스 값을 알 수 있다. 
   - 위 예제에서는 blockIdx.x의 값은 0~N-1의 값을 가진다. 
   - blockIdx의 값이 65535를 초과해서는 안된다. 
(3) 반드시 tid가 실행되는 병렬 갯수 N보다 작은지 점검을 해야한다.  
   - 만일 잘못된 경우, tid가 프로그램이 소유하지 않은 메모리를 가리킬 수 있으며, 잘못된 결과를 주게 될 가능성이 있다. 

 

 

수정 1) 

만일 Thread와 Block을 같이 사용할 경우에는 아래와 같이 수정한다.

...
#define thread_N 5
#deinfe Block_N 2
#define N 10


__global__ void add(int *a, int *b, int *c){
	int tid = threadIdx.x + blockIdx.x * blockDim.x;    
    ...
    
int main(void){
	...
	add<<<Block_N, thread_N>>>(dev_a, dev_b, dev_c);
	...