From Florian, 3 Months ago, written in C++.
Embed
  1. #include "cuda_helper.cuh"
  2.  
  3. __global__
  4. void VectorAddKernel(float *A, float *B, float *C, int N) {
  5.         int i = blockIdx.x * blockDim.x + threadIdx.x;
  6.         if (i < N) {
  7.                 C[i] = A[i] + B[i];
  8.         }
  9. }
  10.  
  11. #define CUDA_MEASURE_START() \
  12.     do { \
  13.                 cudaEvent_t start; \
  14.         cudaEvent_t stop; \
  15.         handleCudaError(cudaEventCreate(&start)); \
  16.         handleCudaError(cudaEventCreate(&stop)); \
  17.         handleCudaError(cudaEventRecord(start));
  18.  
  19. #define CUDA_MEASURE_END(NAME) \
  20.                 handleCudaError(cudaEventRecord(stop)); \
  21.                 handleCudaError(cudaEventSynchronize(stop)); \
  22.                 handleCudaError(cudaEventElapsedTime(&NAME, start, stop)); \
  23.                 handleCudaError(cudaEventDestroy(start)); \
  24.                 handleCudaError(cudaEventDestroy(stop)); \
  25.         } while (0);
  26.  
  27. void cudaVectorAdd(float *a, float *b, float *c, int numElements, int repetitions, bool warmup) {
  28.         float hostToDevice = 0;
  29.         float compute = 0;
  30.         float deviceToHost = 0;
  31.         clock_t start = clock();
  32.  
  33.         int blockSize = 1024;  // FIXME get dynamically?
  34.         int gridSize = (numElements + blockSize - 1) / blockSize;
  35.  
  36.         size_t size = numElements * sizeof(float);
  37.         float *d_A, *d_B, *d_C;
  38.  
  39.         cudaMalloc(&d_A, size);
  40.         cudaMalloc(&d_B, size);
  41.         cudaMalloc(&d_C, size);
  42.  
  43.         CUDA_MEASURE_START();
  44.         cudaMemcpy(d_A, a, size, cudaMemcpyHostToDevice);
  45.         cudaMemcpy(d_B, b, size, cudaMemcpyHostToDevice);
  46.         CUDA_MEASURE_END(hostToDevice);
  47.  
  48.         CUDA_MEASURE_START();
  49.         for (int i = 0; i < repetitions; i++)
  50.         {
  51.                 VectorAddKernel<<<gridSize, blockSize>>>(d_A, d_B, d_C, numElements);
  52.         }
  53.         CUDA_MEASURE_END(compute);
  54.  
  55.         CUDA_MEASURE_START();
  56.         cudaMemcpy(c, d_C, size, cudaMemcpyDeviceToHost);
  57.         CUDA_MEASURE_END(deviceToHost);
  58.  
  59.         cudaFree(d_A);
  60.         cudaFree(d_B);
  61.         cudaFree(d_C);
  62.  
  63.         float total = float(clock() - start) / (CLOCKS_PER_SEC * repetitions);
  64.  
  65.         if (!warmup)
  66.         {
  67.                 printf("CUDA: %.3lf seconds\n", total);
  68.                 printf("CUDA: Copy input to device: %.3lf seconds\n", hostToDevice / (1000 * repetitions));
  69.                 printf("CUDA: Compute time: %.3lf seconds\n", compute / (1000 * repetitions));
  70.                 printf("CUDA: Copy output to host: %.3lf seconds\n", deviceToHost / (1000 * repetitions));
  71.         }
  72. }
  73.  
  74. void fillRandomArray(float *a, int numElements) {
  75.         for (int i = 0; i < numElements; i++) {
  76.                 a[i] = rand() / (float)RAND_MAX;
  77.         }
  78. }
  79.  
  80. void verifyResults(float *a, float *b, float *c, int numElements) {
  81.         for (int i = 0; i < numElements; i++) {
  82.         if (fabs(a[i] + b[i] - c[i]) > 1e-5) {
  83.             fprintf(stderr, "Result verification failed at element %d!\n", i);
  84.             exit(EXIT_FAILURE);
  85.         }
  86.     }
  87. }
  88.  
  89. void sequentialVectorAdd(float *a, float *b, float *c, int numElements) {
  90.         clock_t start = clock();
  91.  
  92.         for (int i = 0; i < numElements; i++) {
  93.                 c[i] = a[i] + b[i];
  94.         }
  95.  
  96.         float diff = float(clock() - start) / CLOCKS_PER_SEC;
  97.         printf("Sequential: %.3lf seconds\n", diff);
  98. }
  99.  
  100. int main() {
  101.         int N = 10000000;
  102.         size_t size = N * sizeof(float);
  103.  
  104.         float *h_a = (float *)malloc(size);
  105.         handleAllocationError(h_a);
  106.         fillRandomArray(h_a, N);
  107.        
  108.         float *h_b = (float *)malloc(size);
  109.         handleAllocationError(h_b);
  110.         fillRandomArray(h_b, N);
  111.        
  112.         float *h_c = (float *)malloc(size);
  113.         handleAllocationError(h_c);
  114.  
  115.         cudaVectorAdd(h_a, h_b, h_c, N, 10, true);
  116.         verifyResults(h_a, h_b, h_c, N);
  117.         cudaVectorAdd(h_a, h_b, h_c, N, 100, false);
  118.  
  119.         sequentialVectorAdd(h_a, h_b, h_c, N);
  120.  
  121.         free(h_a);
  122.         free(h_b);
  123.         free(h_c);
  124.  
  125.         return 0;
  126. }
  127.