Escribí un breve programa de CUDA que usa elbiblioteca CUB altamente optimizada para demostrar que un núcleo de un antiguo procesador Intel Q6600 de cuatro núcleos (los cuatro supuestamente son capaces de ~ 30 GFLOPS / seg) puede hacer un escaneo inclusivo (o suma acumulativa / prefijo si lo prefiere) en 100,000 elementos más rápido que un Nvidia 750 Ti (supuestamente capaz de 1306 GFLOPS / seg de precisión simple). ¿Por qué es este el caso?

El código fuente es:

#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;
}

Y el resultado cuando lo ejecuto en mi máquina es:

Media Tiempo de simulación de GPU: 0.0011999999405816197 [sim / seg]

Media Tiempo de simulación de CPU: 0.00059999997029080987 [sim / seg]

Tiempo de GPU / CPU: 2x

Además, aquí están mis indicadores de compilación y salida:

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 ==========