3. CUDA 예제 - 초급

지금까지 CUDA 에서 데이터 처리하는 방법에 대해 알아보았다. 이번 시간에는 간단한 CUDA 예제를 통해 자세하게 알아보도록 하자. 앞서 GPU에서 데이터 병렬처리를 다음과 같은 순서로 진행한다고 하였다.

CUDA에서 데이터 병렬 처리

1. 입력과 출력에 사용할 데이터를 PC 메모리에 할당

2. 입력과 출력에 사용할 데이터를 그래픽 메모리에 할당

3. 처리하고자 하는 값을 PC 메모리에 입력

4. PC 메모리에 있는 입력 데이터를 그래픽 메모리로 복사

5. 데이터를 분할하여 GPU로 가져옴

6. 수천개 이상의 스레드를 생성하여 커널 함수로 병렬처리

7. 처리된 결과를 병합

8. PC 메모리에 결과를 전송

9. 그래픽 메모리를 해제

10. PC 메모리를 해제

1. 입력과 출력에 사용할 데이터를 PC 메모리에 할당

const int size = 512 * 65535;
const int BufferSize = size * sizeof(int);

int* a;
int* b;
int* c;

a = (int*)malloc(BufferSize);
b = (int*)malloc(BufferSize);
c = (int*)malloc(BufferSize);

입력에 사용될 변수 a와 b와 출력에 사용될 변수 c를 Host에 할당해준다.

2. 입력과 출력에 사용할 데이터를 그래픽 메모리에 할당

int* d_a;
int* d_b;
int* d_c;

cudaMalloc((void**)&d_a, BufferSize);
cudaMalloc((void**)&d_b, BufferSize);
cudaMalloc((void**)&d_c, BufferSize);

1.에서 선언한 변수와 같은 크기의 변수를 Device에 할당해준다. 이때 d_x 변수가 가리키는 주소는 GPU의 메모리이다.

3. 처리하고자 하는 값을 PC 메모리에 입력

int i = 0;

for (i = 0; i < size; i++) {
    a[i] = i;
    b[i] = i;
    c[i] = 0;
}

처리하고자 하는 데이터를 Host에 입력해준다.

4. PC 메모리에 있는 입력 데이터를 그래픽 메모리로 복사

cudaMemcpy(d_a, a, BufferSize, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, BufferSize, cudaMemcpyHostToDevice);

Host에 입력된 데이터를 cudaMemcpy 함수를 사용하여 Device로 복사해준다. 이때 Host에서 Device로 데이터를 복사하는 것이기 때문에 cudaMemcpyKind로 cudaMemcpyHostToDevice를 사용한다.

5. 데이터를 분할하여 GPU로 가져옴

VectorAdd <<<65535, 512 >>> (d_a, d_b, d_c);

CUDA에서 커널 함수를 호출할때 스레드를 같이 생성한다고 했다. 이때 생성되는 스레드의 수에 따라서 데이터를 분할하여 GPU에서 병렬 처리를 할 수 있다.

6. 수천개 이상의 스레드를 생성하여 커널 함수로 병렬처리

__global__ void VectorAdd(int *a, int *b, int *c) {
	int tid = blockIdx.x * blockDim.x + threadIdx.x;

	c[tid] = a[tid] + b[tid];
}

본 예제에서 사용한 커널 함수이다. 간단하게 3개의 데이터를 받아서 입력데이터 a, b를 더한뒤 출력 데이터 c에 저장하는 함수이다. 이때 스레드 인덱스를 접근하여 병렬프로그래밍을 하는데 앞서 커널 함수를 호출할 때 블록수는 65535, 블록 당 스레드의 수는 512로 호출을 하였다. 이는 입출력에 사용될 데이터의 수가 65535 * 512 이기 때문에 같은 수의 스레드를 생성해 준 것이다. 512는 블록이 가질 수 있는 스레드의 최대값이고 65535는 블록의 최대값이다.

7. 처리된 결과를 병합, 8. PC 메모리에 결과를 전송

cudaMemcpy(c, d_c, BufferSize, cudaMemcpyDeviceToHost);

Device에서 커널 함수를 통해 연산된 결과값을 Host로 복사를 한다. 이때 Device에서 Host로 데이터를 복사하는 것이기 때문에 cudaMemcpyKind로 cudaMemcpyDeviceToHost를 사용한다.

9. 그래픽 메모리를 해제

cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

cudaFree 함수를 사용해 할당된 Device 메모리를 해제한다.

10. PC 메모리를 해제

free(a);
free(b);
free(c);

Host 메모리도 마찬가지로 free함수를 사용해 해제한다.

전체 코드

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

__global__ void VectorAdd(int *a, int *b, int *c) {
	int tid = blockIdx.x * blockDim.x + threadIdx.x;

	c[tid] = a[tid] + b[tid];
}

int main() {
	const int size = 512 * 65535;
	const int BufferSize = size * sizeof(int);

	int* a;
	int* b;
	int* c;

	a = (int*)malloc(BufferSize);
	b = (int*)malloc(BufferSize);
	c = (int*)malloc(BufferSize);

	int* d_a;
	int* d_b;
	int* d_c;

	cudaMalloc((void**)&d_a, BufferSize);
	cudaMalloc((void**)&d_b, BufferSize);
	cudaMalloc((void**)&d_c, BufferSize);

	int i = 0;

	for (i = 0; i < size; i++) {
		a[i] = i;
		b[i] = i;
		c[i] = 0;
	}

	cudaMemcpy(d_a, a, BufferSize, cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, b, BufferSize, cudaMemcpyHostToDevice);

	VectorAdd <<<65535, 512 >>> (d_a, d_b, d_c);

	cudaMemcpy(c, d_c, BufferSize, cudaMemcpyDeviceToHost);

	for (i = 0; i < 5; i++) {
		printf(" Result[%d] : %d\n", i, c[i]);
	}
	printf("......\n");
	for (i = size - 5; i < size; i++) {
		printf(" Result[%d] : %d\n", i, c[i]);
	}

	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);

	free(a);
	free(b);
	free(c);

	return 0;
}