벡터 합을 구하는 호스트 프로그램 #include <stdio.h> #inlcude <stdlib.h> #include <string.h> #define NUM_DATA 1024 int main(void) { int* a, * b, * c; int memSize = sizeof(int) * NUM_DATA a = new int[NUM_DATA]; memset(a, 0, memSize); b = new int[NUM_DATA]; memset(b, 0, memSize); c = new int[NUM_DATA]; memset(c, 0, memSize); for (int i = 0; i < NUM_DATA; i++) { a[i] = rand() % 10; b[i] = rand() % 10; } for (int i = 0; i < NUM_DATA; i++) { c[i] = a[i] + b[i]; } delete[] a; delete[] b; delete[] c; } 벡터 합을 구하는 디바이스 프로그램 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #inlcude <stdlib.h> #include <string.h> #define NUM_DATA 1024 __global__ void vecAdd(int* _a, int* _b, int* _c) { int tID = threadIdx.x; _c[tID] = _a[tID] + _b[tID]; } int main(void) { int* a, * b, * c, * hc; int* da, * db, * dc; int memSize = sizeof(int) * NUM_DATA; printf("%d elements, memSize = %d byte\n", NUM_DATA, memSize); // 호스트 디바이스에 메모리 할당 a = new int[NUM_DATA]; memset(a, 0, memSize); b = new int[NUM_DATA]; memset(b, 0, memSize); c = new int[NUM_DATA]; memset(c, 0, memSize); hc = new int[NUM_DATA]; memset(hc, 0, memSize); // 데이터 값 할당 for (int i = 0; i < NUM_DATA; i++) { a[i] = rand() % 10; b[i] = rand() % 10; } // 호스트에서 벡터 합 계산 (성능 비교용) for (int i = 0; i < NUM_DATA; i++) { hc[i] = a[i] + b[i]; } // 디바이스 메모리를 할당 cudaMalloc(&da, memSize); cudaMemset(da, 0, memSize); cudaMalloc(&db, memSize); cudaMemset(db, 0, memSize); cudaMalloc(&dc, memSize); cudaMemset(dc, 0, memSIze); // 호스트에서 디바이스로 데이터 카피 cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice); // 커널 함수 호출 vecAdd <<<1, NUM_DATA>>>(da, db, dc); // 결과를 디바이스에서 호스트로 카피 cudaMemcpy(c, dc, memSize, cudaMemcpyDevcieToHost); // 디바이스 메모리 해제 cudaFree(da); cudaFree(db); cudaFree(dc); // 결과 체크 bool result = true; for (int i = 0; i < NUM_DATA; i++) { if (hc[i] != c[i]) { printf("[%d] The result is not matched! (%d, %d) \n", i, hc[i], c[i]) result = false } } if (result) printf("GPU works ~") // 호스트 메모리 해제 delete[] a;delete[] b; delete[] c; return 0; } 실행 시간 커널 실행 시간 커널 실행 시간을 측정하기 위해서는 커널 호출 전과 후에 시간을 측정해야 한다. 이를 통해 디바이스의 연산 실행 시간을 알 수 있다. 비동기성과 동기화 호스트는 커널을 호출하면 디바이스에게 커널 실행을 요청하고, 바로 다음 명령으로 넘어간다. 따라서 호스트는 디바이스와 별개로 동시에 작업이 가능하다. 따라서 커널 실행 시간을 정확히 측정하려면, 호스트가 커널 호출 후 바로 다음 작업으로 넘어가는 비동기성을 고려해야 한다. 호스트가 디바이스의 연산 수행이 끝나기까지 기다리게 하여 정확한 시간을 측정할 수 있다. 이를 동기화라고 한다. 주의사항 CUDA API 호출은 큐 구조를 통해 관리된다. 따라서 순차 실행을 한다. 호스트가 디바이스 제어만을 위한 경우, 동기화는 필요없다. 그러나 호스트, 디바이스 모두 연산에 관여하는 경우, 동기화가 필요하다. CUDA API 호출 자체에도 일정 시간의 소모 비용이 있다. 하지만 CUDA 연산의 정확성과 효율성을 모두 잡으려면, 스레드 계층과 스레드 인덱싱 방법이 필요하다. # cudaDeviceSynchronize를 통해 디바이스 연산이 끝날 때까지, 호스트를 기다리게 한다. veeAdd <<<1, NUM_DATA>>>(da, db, dc); cudaDeviceSynchronize(); 데이터 전송 CUDA 프로그램 흐름 호스트에서 디바이스로 데이터 복사 디바이스 연산 디바이스에서 호스트로 데이터 복사 데이터 복사
CUDA 알고리즘의 성능을 판단할 때는, 반드시 데이터 전송 시간을 고려해야 한다. 이 외에도 데이터 변환 등의 추가 작업이 있다면, 이의 소요 시간도 포함해야 한다.