벡터 합을 구하는 호스트 프로그램

#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 프로그램 흐름
    1. 호스트에서 디바이스로 데이터 복사
    2. 디바이스 연산
    3. 디바이스에서 호스트로 데이터 복사
  • 데이터 복사
    CUDA 알고리즘의 성능을 판단할 때는, 반드시 데이터 전송 시간을 고려해야 한다. 이 외에도 데이터 변환 등의 추가 작업이 있다면, 이의 소요 시간도 포함해야 한다.