CUDA C - CPU and GPU execution time using clock() and cudaEvent, is it correct?
c, cuda
Solution
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
Problem
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 ?