1
\$\begingroup\$

I have the default code snippet that comes with the nvidia nsight eclipse V10.2, but I've modified it to do some profiling:

/*
 ============================================================================
 Name        : Project1.cu
 Author      : []
 Version     :
 Copyright   : If you find this, by pure chance, feel free to use the code!
 Description : CUDA compute reciprocals
 ============================================================================
 */

#include <iostream>
#include <numeric>
#include <stdlib.h>
#include <chrono>

static void CheckCudaErrorAux (const char *, unsigned, const char *, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

/**
 * CUDA kernel that computes reciprocal values for a given vector
 */
__global__ void reciprocalKernel(float *data, unsigned vectorSize) {
    unsigned idx = blockIdx.x*blockDim.x+threadIdx.x;
    if (idx < vectorSize)
    {
        data[idx] = 1.0/data[idx];
    }
}

/**
 * Host function that copies the data and launches the work on GPU
 */
void gpuReciprocal(float* gpuData,unsigned size)
{


    //CUDA_CHECK_RETURN(
            //);
    //CUDA_CHECK_RETURN(//);

    static const int BLOCK_SIZE = 512;
    const int blockCount = (size+BLOCK_SIZE-1)/BLOCK_SIZE;
    reciprocalKernel<<<blockCount, BLOCK_SIZE>>> (gpuData, size);

    //CUDA_CHECK_RETURN(
            //);
    //CUDA_CHECK_RETURN(
            //);
}

float* copyMemoryFromDevice(float*& gpuData, unsigned size){
    float *rc = new float[size];
    cudaMemcpy(rc, gpuData, sizeof(float)*size, cudaMemcpyDeviceToHost);
    return rc;
}

void copyMemoryToDevice(float*& gpuData, float* data, unsigned size){
    cudaMemcpy(gpuData, data, sizeof(float)*size, cudaMemcpyHostToDevice);
}

void allocateGPUMemory(float*& gpuData, unsigned size){
    cudaMalloc((void **)&gpuData, sizeof(float)*size);
}

void unAllocateGPUMemory(float*& gpuData){
    cudaFree(gpuData);
}

float* cpuReciprocal(float*& data, unsigned size)
{
    float *rc = new float[size];
    for (unsigned cnt = 0; cnt < size; ++cnt) rc[cnt] = 1.0/data[cnt];
    return rc;
}


void initialize(float *data, unsigned size)
{
    for (unsigned i = 0; i < size; ++i)
        data[i] = .1*(i+1);
}

uint64_t getCurrentTimeMs()
{
    return std::chrono::system_clock::now().time_since_epoch().count();
}

std::string getCommaFormattedNumber( uint64_t number ){
    std::string toreturn = std::to_string(number);
    int placeholder = 3;
    while(placeholder < toreturn.size()){
        toreturn.insert(toreturn.size()-placeholder, ",");
        placeholder += 4;
    }
    return toreturn;
}

int main(void)
{
    uint64_t startCPU,endCPU,CPUTime
    ,startGPU,endGPU,GPUTime
    ,startAcum,endAcum,ACUMTime
    ,startMalloc,endMalloc,mallocTime
    ,startCopy,endCopy,copyTime
    ,startCompute,endCompute,computeTime
    ,percentDiff;
    std::string CPUTimeFormatted,GPUTimeFormatted,ACUMTimeFormatted,mallocTimeFormatted,copyTimeFormatted,computeTimeFormatted,percentDiffFormatted;

    static const int WORK_SIZE = 292565536;
    float* data = new float[WORK_SIZE];
    initialize (data, WORK_SIZE);
    cudaSetDevice(0);

    //CPU computation phase
    std::cout<<"Started CPU"<<std::endl;
    startCPU = getCurrentTimeMs();
    float *recCpu = cpuReciprocal(data, WORK_SIZE);
    endCPU = getCurrentTimeMs();

    //GPU computation phase
    std::cout<<"Started GPU"<<std::endl;
    startGPU = getCurrentTimeMs();

    startMalloc = getCurrentTimeMs();
    float* gpuData;
    allocateGPUMemory(gpuData, WORK_SIZE);
    endMalloc = getCurrentTimeMs();

    startCopy = getCurrentTimeMs();
    copyMemoryToDevice(gpuData, *&data, WORK_SIZE);
    endCopy = getCurrentTimeMs();

    startCompute = getCurrentTimeMs();
    gpuReciprocal(gpuData, WORK_SIZE);
    endCompute = getCurrentTimeMs();

    startCopy += getCurrentTimeMs();
    float *recGpu = copyMemoryFromDevice(gpuData, WORK_SIZE);
    endCopy += getCurrentTimeMs();

    startMalloc += getCurrentTimeMs();
    unAllocateGPUMemory(gpuData);
    endMalloc += getCurrentTimeMs();

    endGPU = getCurrentTimeMs();

    //Accumulation phase
    std::cout<<"Started acum"<<std::endl;
    startAcum = getCurrentTimeMs();
    float cpuSum = std::accumulate (recCpu, recCpu+WORK_SIZE, 0.0);
    float gpuSum = std::accumulate (recGpu, recGpu+WORK_SIZE, 0.0);
    endAcum = getCurrentTimeMs();

    //Results calculation
    CPUTime = (endCPU - startCPU)/1000;
    GPUTime = (endGPU - startGPU)/1000;
    ACUMTime = (endAcum - startAcum)/1000;
    mallocTime = (endMalloc - startMalloc)/1000;
    copyTime = (endCopy - startCopy)/1000;
    computeTime = (endCompute - startCompute)/1000;
    percentDiff = (CPUTime*100)/GPUTime;
    CPUTimeFormatted = getCommaFormattedNumber(CPUTime);
    GPUTimeFormatted = getCommaFormattedNumber(GPUTime);
    ACUMTimeFormatted = getCommaFormattedNumber(ACUMTime);
    mallocTimeFormatted = getCommaFormattedNumber(mallocTime);
    copyTimeFormatted = getCommaFormattedNumber(copyTime);
    computeTimeFormatted = getCommaFormattedNumber(computeTime);
    percentDiffFormatted = getCommaFormattedNumber(percentDiff);

    /* Verify the results */
    std::cout<<"gpuSum = "<<gpuSum<< " cpuSum = " <<cpuSum<<std::endl;
    std::cout<<"CPU: "<<CPUTimeFormatted<<"us "
            <<std::endl<<"GPU: "<<GPUTimeFormatted<<"us "
            <<std::endl<<" MALLOC: "<<mallocTimeFormatted<<"us "
            <<std::endl<<" COPY: "<<copyTimeFormatted<<"us "
            <<std::endl<<" COMPUTE: "<<computeTimeFormatted<<"us "
            <<std::endl<<"ACUM: "<<ACUMTimeFormatted<<"us"
            <<std::endl;
    std::cout<<"The GPU is "<<percentDiffFormatted<<"% faster than the CPU.";

    /* Free memory */
    delete[] data;
    delete[] recCpu;
    delete[] recGpu;

    return 0;
}

/**
 * Check the return value of the CUDA runtime API call and exit
 * the application if the call has failed.
 */
static void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err)
{
    if (err == cudaSuccess)
        return;
    std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
    exit (1);
}

And the code produces the following results:

CPU: 1,069,778us
GPU: 562,895us
MALLOC: 43,535us
COPY: 519,350us
COMPUTE: 7us
ACUM: 1,630,715us
The GPU is 190% faster than the CPU.

As one can see, the copying from RAM to GPU and back costs the most time, and the malloc / demalloc also takes a large amount of time. This seems very strange, as the GPU I'm using is on a 8X PCIE lane. I also have a NVidia Jetson Xavier, and when I run this code on that system, I get the following output:

CPU: 774,552us
GPU: 446,498us
MALLOC: 103,619us
COPY: 342,860us
COMPUTE: 14us
ACUM: 967,493us
The GPU is 173% faster than the CPU.logout

As one can see, the compute call is by far the least intensive call, and the malloc/copy calls are almost ludicrous in terms of real time operation.

How can I optimize this snippet to have faster copy and or malloc times?

\$\endgroup\$
3
  • 1
    \$\begingroup\$ The profiling code is wrong, for a common reason: no synchronization (meaning that what was measured is kernel launch time, not kernel execution time). That might invalidate this whole question.. \$\endgroup\$ – harold May 28 '20 at 21:15
  • \$\begingroup\$ @harold even if the execution is not synchronized, the malloc and demalloc times are still valid and are still large in comparison \$\endgroup\$ – tuskiomi May 28 '20 at 21:28
  • \$\begingroup\$ Firstly, as harold says, you should use CUDA events for measuring performance: How to Implement Performance Metrics in CUDA C/C++. Nonetheless, your observation is valid – the memory (de)allocation and transfer times are considerable compared to the computation time. The reason is that your computation is simply too trivial to take any substantial amount of time. You can improve CPU <-> GPU transfer speed using pinned memory; read devblogs.nvidia.com/how-optimize-data-transfers-cuda-cc. \$\endgroup\$ – kyrill May 30 '20 at 0:14

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.