I wrote a short CUDA program that uses the highly-optimized CUB library to demonstrate that one core from an old, quad-core Intel Q6600 processor (all four are supposedly capable of ~30 GFLOPS/sec) can do an inclusive scan (or cumulative/prefix sum if you rather) on 100,000 elements faster than an Nvidia 750 Ti (supposedly capable of 1306 GFLOPS/sec of single precision). Why is this the case?
The source code is:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cub/cub.cuh>
#include <stdio.h>
#include <time.h>
#include <algorithm>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
if (code != cudaSuccess)
{
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
void fillArrayWithRandom(float* inputArray, int inputN)
{
for (int i = 0; i < inputN; i++)
{
inputArray[i] = (float)rand() / float(RAND_MAX);
}
}
void inclusiveSum_CPU(float *inputArray, float *inputSummedArray, int inputN)
{
for (int i = 0; i < inputN; i++)
{
if (i > 0)
{
inputSummedArray[i] = inputSummedArray[i - 1] + inputArray[i];
}
else
{
inputSummedArray[i] = inputArray[i];
}
}
}
int main()
{
int N = 100000; //1 hundred thousand elements
float numSimulations = 10000;
//Make Host Arrays
float* testArray_CPU = (float *)malloc(sizeof(float)*N);
fillArrayWithRandom(testArray_CPU, N);
float* testArrayOutput_CPU = (float *)malloc(sizeof(float)*N);
//Make GPU Arrays
float* testArray_GPU;
gpuErrchk(cudaMalloc(&testArray_GPU, N*sizeof(float)));
gpuErrchk(cudaMemcpy(testArray_GPU, testArray_CPU, N*sizeof(float), cudaMemcpyHostToDevice));
float* testArrayOutput_GPU;
gpuErrchk(cudaMalloc(&testArrayOutput_GPU, N*sizeof(float)));
//Initiate the benchmark variables
clock_t begin_CPU, end_CPU;
float time_spent_GPU, time_spent_CPU;
//GPU prep
void *d_temp_storage = NULL;
size_t temp_storage_bytes = 0;
cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, testArray_GPU, testArrayOutput_GPU, N);
gpuErrchk(cudaMalloc(&d_temp_storage, temp_storage_bytes));
//GPU Timing
cudaEvent_t start, stop;
gpuErrchk(cudaEventCreate(&start));
gpuErrchk(cudaEventCreate(&stop));
gpuErrchk(cudaEventRecord(start, 0));
for (int i = 0; i < numSimulations; i++)
{
cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, testArray_GPU, testArrayOutput_GPU, N);
}
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaEventRecord(stop, 0));
gpuErrchk(cudaEventSynchronize(stop));
gpuErrchk(cudaEventElapsedTime(&time_spent_GPU, start, stop));
cudaError_t error = cudaGetLastError();
if (error != cudaSuccess)
{
printf("CUDA error: %s\n", cudaGetErrorString(error));
exit(-1);
}
time_spent_GPU = (float)(time_spent_GPU / 1000);
float avg_GPU = time_spent_GPU / numSimulations;
printf("Avg. GPU Simulation Time: %.17g [sim/sec]\n", avg_GPU);
//CPU Timing
begin_CPU = clock();
for (int i = 0; i < numSimulations; i++)
{
inclusiveSum_CPU(testArray_CPU, testArrayOutput_CPU, N);
}
end_CPU = clock();
time_spent_CPU = (float)(end_CPU - begin_CPU) / CLOCKS_PER_SEC;
float avg_CPU = time_spent_CPU / numSimulations;
printf("Avg. CPU Simulation Time: %.17g [sim/sec]\n", avg_CPU);
printf("GPU/CPU Timing:%.17gx \n", avg_GPU / avg_CPU);
return 0;
}
And the output when I run it on my machine is:
Avg. GPU Simulation Time: 0.0011999999405816197 [sim/sec]
Avg. CPU Simulation Time: 0.00059999997029080987 [sim/sec]
GPU/CPU Timing:2x
Also, here are my compiling flags and output:
1>------ Build started: Project: speedTest, Configuration: Debug Win32 ------
1> Compiling CUDA source file kernel.cu...
1>
1> C:\Users\Owner\Documents\Visual Studio 2013\Projects\speedTest\speedTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\bin\nvcc.exe" -gencode=arch=compute_50,code=\"sm_50,compute_50\" --use-local-env --cl-version 2013 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 12.0\VC\bin" -rdc=true -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\include" -G --keep-dir Debug -maxrregcount=0 --machine 32 --compile -cudart static -g -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -o Debug\kernel.cu.obj "C:\Users\Owner\Documents\Visual Studio 2013\Projects\speedTest\speedTest\kernel.cu"
1> kernel.cu
1>
1> C:\Users\Owner\Documents\Visual Studio 2013\Projects\speedTest\speedTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\bin\nvcc.exe" -dlink -o Debug\speedTest.device-link.obj -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\lib\Win32" cudart.lib kernel32.lib user32.lib gdi32.lib winspool.lib comdlg32.lib advapi32.lib shell32.lib ole32.lib oleaut32.lib uuid.lib odbc32.lib odbccp32.lib -gencode=arch=compute_50,code=sm_50 -G --machine 32 Debug\kernel.cu.obj
1> cudart.lib
1> kernel32.lib
1> user32.lib
1> gdi32.lib
1> winspool.lib
1> comdlg32.lib
1> advapi32.lib
1> shell32.lib
1> ole32.lib
1> oleaut32.lib
1> uuid.lib
1> odbc32.lib
1> odbccp32.lib
1> kernel.cu.obj
1> speedTest.vcxproj -> C:\Users\Owner\Documents\Visual Studio 2013\Projects\speedTest\Debug\speedTest.exe
1> copy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\bin\cudart*.dll" "C:\Users\Owner\Documents\Visual Studio 2013\Projects\speedTest\Debug\"
1> C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\bin\cudart32_65.dll
1> C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\bin\cudart64_65.dll
1> 2 file(s) copied.
========== Build: 1 succeeded, 0 failed, 0 up-to-date, 0 skipped ==========
