I want to use clock() to compare different kernel implementations. I tried to implement it in a simple SAXPY example but it leads to zero clock cycles, which is very unlikely.
I already found some examples on how to implement the clock(). here and here. But somehow the transfer to my code does not work.
Here is the code that I am using:
/* SAXPY code example from https://devblogs.nvidia.com/parallelforall/easy-introduction-cuda-c-and-c/ */
#include <stdio.h>
// The declaration specifier __global__ defines a kernel. This code
// will be copied to the device and will be executed there in parallel
__global__
void saxpy(int n, float a, float *x, float *y, int *kernel_clock)
{
// The indexing of the single threads is done with the following
// code line
int i = blockIdx.x*blockDim.x + threadIdx.x;
clock_t start = clock();
// Each thread is executing just one position of the arrays
if (i < n) y[i] = a*x[i] + y[i];
clock_t stop = clock();
kernel_clock[i] = (int) (stop-start);
}
int main(void)
{
// Clock cycles of threads
int *kernel_clock;
int *d_kernel_clock;
// Creating a huge number
int N = 1<<20;
float *x, *y, *d_x, *d_y;
// Allocate an array on the *host* of the size of N
x = (float*)malloc(N*sizeof(float));
y = (float*)malloc(N*sizeof(float));
kernel_clock = (int*)malloc(N*sizeof(int));
// Allocate an array on the *device* of the size of N
cudaMalloc(&d_x, N*sizeof(float));
cudaMalloc(&d_y, N*sizeof(float));
cudaMalloc(&d_kernel_clock, N*sizeof(int));
// Filling the array of the host
for (int i = 0; i < N; i++) {
x[i] = 1.0f;
y[i] = 2.0f;
}
// Copy the host array to the device array
cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_kernel_clock, kernel_clock, N*sizeof(int), cudaMemcpyHostToDevice);
// Perform SAXPY on 1M elements. The triple chevrons dedicates how
// the threads are grouped on the device
saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y, d_kernel_clock);
cudaDeviceSynchronize();
// Copy the result from the device to the host
cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);
cudaMemcpy(kernel_clock, d_kernel_clock, N*sizeof(int), cudaMemcpyDeviceToHost);
// Calculate average clock time
float average_clock = 0;
for (int i = 0; i < N; i++) {
average_clock += (float) (kernel_clock[i]);
}
average_clock /= N;
// Display the time to the screen
printf ("Kernel clock cycles: %.4f\n", average_clock);
// Free the memory on the host and device
free(x);
free(y);
free(kernel_clock);
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_kernel_clock);
}
This code example leads to:
Kernel clock cycles: 0.0000
I am not sure what I am doing wrong. So my question is: How do I actually get a reasonable result?
