다음은 두 Array의 합을 GPU로 구하는 간단한 코드이다.
#include <stdio.h> #define N 1000 __global__ void add(int *a,int *b,int *c){ /* (1) */ int tid=blockIdx.x; 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; /* (2) */ 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; } /* (3) */ cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice); cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice); /* (4) */ add<<<N,1>>>(dev_a,dev_b,dev_c); /* (5) */ 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]); /* (6) */ cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; } |
(1) blockIdx.x 는 커널 복사본의 Index 이다. (4)에서 kernel<<N,M>>{param1,param2,...} 에서 N은 커널 복사본이고, M은 하나의 블록안에 실행할 쓰레드 개수이다. (http://alnova2.tistory.com/818) 이 예제에서는 N개 만큼의 커널 복사본을 실행한다.
(2) GPU에 메모리를 할당한다.
(3) CPU 메모리의 값을 GPU 메모리로 옮긴다.
(4) Kernel 을 시작한다.
(5) GPU 메모리값을 를 CPU 메모리로 옮긴다
(6) GPU 메모리를 해제한다.
이 예제는 N개 만큼의 add 코드를 병렬로 작동 시킨 예이다.