CUDA C - CPU and GPU execution time using clock() and cudaEvent, is it correct?

Question

I wrote a program to add Two 2D array to check the performance of CPU and GPU. I used clock() function to measure the CPU execution and cudaEvent to measure the kernal execution time in GPU. Since I was learning CUDA under Udacity, i tried to execute the program on their server and found the result as,

 Output:
 GPU: 0.001984 ms
 CPU : 30.000000 ms

And now down to my real question, I find these result to be amazingly fast on the GPU and now I'm a bit skeptical on whether these results are accurate or have i done any mistake in my program ?

Here's my program :

 #include "stdio.h"
 #include<time.h>
 #define COLUMNS 900
 #define ROWS 900
 long a[ROWS][COLUMNS], b[ROWS][COLUMNS], c[ROWS][COLUMNS],d[ROWS][COLUMNS];
__global__ void add(long *a, long *b, long *c,long *d)
{
 int x = blockIdx.x;
 int y = blockIdx.y;
 int i = (COLUMNS*y) + x;
 c[i] = a[i] + b[i];
 a[i]=d[i];
}

int main()
{
  long *dev_a, *dev_b, *dev_c,*dev_d;
  float ms;
  clock_t startc, end;
  double cpu_time_used;
  cudaEvent_t start,stop;


 cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_d, ROWS*COLUMNS*sizeof(int));

 startc = clock();
 for (long y = 0; y < ROWS; y++) // Fill Arrays
 for (long x = 0; x < COLUMNS; x++)
 {
     a[y][x] = x;
     b[y][x] = y;
     d[y][x]=rand()%4;
     c[y][x]=a[y][x]+b[y][x];
 }
 end = clock();

cpu_time_used = ((double) (end - startc)) / CLOCKS_PER_SEC;
cpu_time_used*=1000;


cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_d, d, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);


cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0); 
cudaEventRecord(stop, 0);


add<<<dim3(1024,1024),dim3(128,128)>>>(dev_a, dev_b, dev_c,dev_d);

cudaEventSynchronize(stop);
cudaEventElapsedTime(&ms, start, stop);
cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int),cudaMemcpyDeviceToHost);
cudaEventDestroy(start);
cudaEventDestroy(stop);




printf("GPU: %f ms",ms);
printf("\n CPU : %f ms",cpu_time_used);

 return 0;
}

Thank you all for the answer provided to my query and here are the changes I made to the code and the updated result,

Updated Code:

#include "stdio.h"
#include <time.h>
#include <sys/time.h>
#include <unistd.h>
#define COLUMNS 500
#define ROWS 500
long a[ROWS][COLUMNS], b[ROWS][COLUMNS], c[ROWS][COLUMNS],d[ROWS][COLUMNS];



__global__ void add(long *a, long *b, long *c,long *d)
{
 int x = blockIdx.x;
 int y = blockIdx.y;
 int i = (COLUMNS*y) + x;
 c[i] = a[i] + b[i];
 a[i]=d[i];
}
int main()
{
 long *dev_a, *dev_b, *dev_c,*dev_d;
 struct timeval startc, end;
 float ms;
 long mtime, seconds, useconds;
 //   clock_t startc, end;
 //  double cpu_time_used;
 long ns;
 cudaEvent_t start,stop;


 cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_d, ROWS*COLUMNS*sizeof(int));

 gettimeofday(&startc, NULL);
 for (long y = 0; y < ROWS; y++) // Fill Arrays
 for (long x = 0; x < COLUMNS; x++)
 {
  a[y][x] = x;
  b[y][x] = y;
  d[y][x]=rand()%4;
  c[y][x]=a[y][x]+b[y][x];
 }
  gettimeofday(&end, NULL);

 seconds  = end.tv_sec  - startc.tv_sec;
 useconds = end.tv_usec - startc.tv_usec;
 mtime = ((seconds) * 1000 + useconds/1000.0) + 0.5;


for (long y = ROWS-1; y < ROWS; y++) // Output Arrays
 {
 for (long x = COLUMNS-1; x < COLUMNS; x++)
 {
    // printf("\n[%ld][%ld]=%ld ",y,x,c[y][x]);
   //   printf("[%d][%d]=%d ",y,x,d[y][x]);
 }
 printf("\n");
 }



cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_d, d, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);


cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0); 



add<<<dim3(1024,1024),dim3(128,128)>>>(dev_a, dev_b, dev_c,dev_d);

cudaThreadSynchronize();
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&ms, start, stop);

cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int),cudaMemcpyDeviceToHost);
cudaEventDestroy(start);
cudaEventDestroy(stop);



//cpu_time_used = ((double) (end - start)) / CLOCKS_PER_SEC;
printf("GPU: %f ms",ms);
printf("\n CPU : %ld ms",mtime);
 for (long y = ROWS-1; y < ROWS; y++) // Output Arrays
 {
     for (long x = COLUMNS-1; x < COLUMNS; x++)
     {
      //   printf("\n[%ld][%ld]=%ld ",y,x,c[y][x]);
      //   printf("[%d][%d]=%d ",y,x,d[y][x]);
     }
     printf("\n");
 }
 return 0;
}

Output:

GPU: 0.011040 ms
CPU : 9 ms

Now can i safely tell whether its correct ?

The answers to [Timing CUDA operations](http://stackoverflow.com/questions/7876624/timing-cuda-operations) may be of your interest. — Vitality, Feb 24 '15 at 21:06

score 3 · Accepted Answer · edited May 23 '17 at 12:17

You are correct in thinking that is too much of a speedup, the timing for the CPU is too long. Use this method to time the CPU C++ obtaining milliseconds time on Linux -- clock() doesn't seem to work properly also you may have to move the cudaEventRecord(stop, 0); to after the kernel.

I see 5 read and writes in your kernel. Taking 5*4Bytes*500*500/(1024^3*0.009) You're getting about 0.517 GB/s out of your memory, this is a small fraction of available. I would say your CPU version needs some work. In contrast your GPU is at 5*4Bytes*500*500/(1024^3*0.01104e-3) is about 421GB/s. I would say you're not quite there.

So, so many errors....

#include "stdio.h"
#include <time.h>
#include <sys/time.h>
#include <unistd.h>
#include <cuda.h>
#include <cuda_runtime.h>

#define COLUMNS 500
#define ROWS 500
long a[ROWS*COLUMNS], b[ROWS*COLUMNS], c[ROWS*COLUMNS],d[ROWS*COLUMNS];



__global__ void add(long *a, long *b, long *c,long *d)
{
 int x = blockIdx.x;
 int y = blockIdx.y;
 int i = (COLUMNS*y) + x;
 c[i] = a[i] + b[i];
 a[i]=d[i];
}
int main()
{
 long *dev_a, *dev_b, *dev_c,*dev_d;
 struct timeval startc, end;
 float ms;
 long seconds, useconds;
 double mtime;
 cudaEvent_t start,stop;


 for(int i=0; i<ROWS*COLUMNS; i++)
     d[i]=rand()%4;

 for(int i=0; i<ROWS; i++){
     for(int j=0; j<COLUMNS; j++){
         a[i*COLUMNS+j]=j;
         b[i*COLUMNS+j]=i;
     }
 }

 cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int));
 cudaMalloc((void **) &dev_d, ROWS*COLUMNS*sizeof(int));



 gettimeofday(&startc, NULL);
 for (long i = 0; i < ROWS*COLUMNS; i++){ // Fill Arrays
     c[i]=a[i]+b[i];
     a[i]=d[i];
 }
  gettimeofday(&end, NULL);

 seconds  = end.tv_sec  - startc.tv_sec;
 useconds = end.tv_usec - startc.tv_usec;
 mtime = useconds;
 mtime/=1000;
 mtime+=seconds*1000;

for (long y = ROWS-1; y < ROWS; y++) // Output Arrays
 {
 for (long x = COLUMNS-1; x < COLUMNS; x++)
 {
    // printf("\n[%ld][%ld]=%ld ",y,x,c[y][x]);
   //   printf("[%d][%d]=%d ",y,x,d[y][x]);
 }
 printf("\n");
 }



cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_d, d, ROWS*COLUMNS*sizeof(int),
cudaMemcpyHostToDevice);


cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);



add<<<dim3(1024,1024),dim3(128,128)>>>(dev_a, dev_b, dev_c,dev_d);



cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&ms, start, stop);

cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int),cudaMemcpyDeviceToHost);
cudaEventDestroy(start);
cudaEventDestroy(stop);

printf("GPUassert: %s\n", cudaGetErrorString(cudaGetLastError()));

//cpu_time_used = ((double) (end - start)) / CLOCKS_PER_SEC;
double memXFers=5*4*COLUMNS*ROWS;
memXFers/=1024*1024*1024;


printf("GPU: %f ms bandwidth %g GB/s",ms, memXFers/(ms/1000.0));
printf("\n CPU : %g ms bandwidth %g GB/s",mtime, memXFers/(mtime/1000.0));
 for (long y = ROWS-1; y < ROWS; y++) // Output Arrays
 {
     for (long x = COLUMNS-1; x < COLUMNS; x++)
     {
      //   printf("\n[%ld][%ld]=%ld ",y,x,c[y][x]);
      //   printf("[%d][%d]=%d ",y,x,d[y][x]);
     }
     printf("\n");
 }

 return 0;
}

My current result by the way (obviously not correct)...

GPU: 0.001792 ms bandwidth 2598.56 GB/s
CPU : 0.567 ms bandwidth 8.21272 GB/s

You *definitely* have to move the `cudaEventRecord(stop, 0);` to after the kernel. — ArchaeaSoftware, Feb 25 '15 at 08:35
First of all thank you for the awesome response, yes i have moved the `cudaEventRecord(stop, 0);` after the kernal execution. — Avinash, Feb 25 '15 at 15:00
and also thanks for the timer function, i tried to implement it in my code and found the result to be promiising. — Avinash, Feb 25 '15 at 15:37
One more question regarding timers, since the timer I used is only limited to linux, is there any similar timer for windows platform? — Avinash, Feb 25 '15 at 15:59
Two options are provided in http://stackoverflow.com/questions/25615571/how-to-use-sys-time-h-in-window (see answer) — Christian Sarofeen, Feb 25 '15 at 21:00
Accept the answer since the question you asked has been answered. Continue figuring out remaining issues in another question. There are a few hints in the posted code. — Christian Sarofeen, Feb 25 '15 at 21:33

CUDA C - CPU and GPU execution time using clock() and cudaEvent, is it correct?

1 Answers1