Temporización de operaciones CUDA

8 minutos de lectura

avatar de usuario
Tudor

Necesito cronometrar una ejecución del kernel CUDA. La Guía de mejores prácticas dice que podemos usar eventos o funciones de tiempo estándar como clock() en Windows. Mi problema es que usar estas dos funciones me da un resultado totalmente diferente. De hecho, el resultado que dan los eventos parece enorme en comparación con la velocidad real en la práctica.

Para lo que realmente necesito todo esto es para poder predecir el tiempo de ejecución de un cálculo ejecutando primero una versión reducida en un conjunto de datos más pequeño. Desafortunadamente, los resultados de este punto de referencia son totalmente poco realistas, ya sea demasiado optimistas (clock()) o demasiado pesimista (eventos).

  • ¿Sincronizó en la CPU después de iniciar su kernel y antes de cronometrar (finalizar) con el reloj?

    – pQB

    24 de octubre de 2011 a las 13:53

  • ¿Quieres decir si tengo llamadas cudaThreadSynchronize() antes y después? Sí.

    – Tudor

    24/10/2011 a las 13:55


  • Sí, eso es lo que quise decir.

    – pQB

    24 de octubre de 2011 a las 14:01

  • Por cierto. El tiempo de CUDA se devuelve en ms (microsegundos si usa el perfilador visual). Por si acaso.

    – pQB

    24/10/2011 a las 14:30


avatar de usuario
fbielejec

Podrías hacer algo a lo largo de las líneas de:

#include <sys/time.h>

struct timeval t1, t2;

gettimeofday(&t1, 0);

kernel_call<<<dimGrid, dimBlock, 0>>>();

HANDLE_ERROR(cudaThreadSynchronize();)

gettimeofday(&t2, 0);

double time = (1000000.0*(t2.tv_sec-t1.tv_sec) + t2.tv_usec-t1.tv_usec)/1000.0;

printf("Time to generate:  %3.1f ms \n", time);

o:

float time;
cudaEvent_t start, stop;

HANDLE_ERROR( cudaEventCreate(&start) );
HANDLE_ERROR( cudaEventCreate(&stop) );
HANDLE_ERROR( cudaEventRecord(start, 0) );

kernel_call<<<dimGrid, dimBlock, 0>>>();

HANDLE_ERROR( cudaEventRecord(stop, 0) );
HANDLE_ERROR( cudaEventSynchronize(stop) );
HANDLE_ERROR( cudaEventElapsedTime(&time, start, stop) );

printf("Time to generate:  %3.1f ms \n", time);

  • @Programmer Es una función o macro que no definió que maneja los errores devueltos por las llamadas a la función cuda. Debería hacer el manejo de errores, pero podría haberse omitido aquí por simplicidad.

    – jmsu

    25 de octubre de 2011 a las 10:21

  • @ Programador, sí, exactamente, hay algunas macros útiles para el manejo de errores en el SDK

    – fbielejec

    25 de octubre de 2011 a las 13:13


  • @fbielejec, las funciones de manejo de errores en el SDK están ahí solo para hacer que los ejemplos sean lo más simples posible para la educación. ¡Llamar a exit() cuando encuentra un error no es la mejor manera de manejar un error!

    – Tomás

    26 de octubre de 2011 a las 10:56

  • Tenga en cuenta que 1e6 us = 1 s, por lo que en el primer ejemplo, el tiempo está en segundos, no en ms.

    – Kknd

    4 de diciembre de 2013 a las 0:17

  • Con respecto a HANDLE_ERROR, consulte stackoverflow.com/q/14038589/2778484 y observe helper_cuda.h en los ejemplos de CUDA, que tiene una macro llamada getLastCudaError.

    – chappjc

    5 de septiembre de 2014 a las 1:37


Ya se ha dado una respuesta satisfactoria a su pregunta.

He construido clases para cronometrar C/C++ así como operaciones CUDA y quiero compartir con otros con la esperanza de que puedan ser útiles para los próximos usuarios. Solo tendrás que agregar el 4 archivos informados a continuación para su proyecto y #include los dos archivos de encabezado como

// --- Timing includes
#include "TimingCPU.h"
#include "TimingGPU.cuh"

Las dos clases se pueden utilizar de la siguiente manera.

Sección de CPU de temporización

TimingCPU timer_CPU;

timer_CPU.StartCounter();
CPU perations to be timed
std::cout << "CPU Timing = " << timer_CPU.GetCounter() << " ms" << std::endl;

Sección GPU de temporización

TimingGPU timer_GPU;
timer_GPU.StartCounter();
GPU perations to be timed
std::cout << "GPU Timing = " << timer_GPU.GetCounter() << " ms" << std::endl;

En ambos casos, el tiempo es en milisegundos. Además, las dos clases se pueden usar en Linux o Windows.

Aquí están los 4 archivos:

TimingCPU.cpp

/**************/
/* TIMING CPU */
/**************/

#include "TimingCPU.h"

#ifdef __linux__

    #include <sys/time.h>
    #include <stdio.h>

    TimingCPU::TimingCPU(): cur_time_(0) { StartCounter(); }

    TimingCPU::~TimingCPU() { }

    void TimingCPU::StartCounter()
    {
        struct timeval time;
        if(gettimeofday( &time, 0 )) return;
        cur_time_ = 1000000 * time.tv_sec + time.tv_usec;
    }

    double TimingCPU::GetCounter()
    {
        struct timeval time;
        if(gettimeofday( &time, 0 )) return -1;

        long cur_time = 1000000 * time.tv_sec + time.tv_usec;
        double sec = (cur_time - cur_time_) / 1000000.0;
        if(sec < 0) sec += 86400;
        cur_time_ = cur_time;

        return 1000.*sec;
    }

#elif _WIN32 || _WIN64
    #include <windows.h>
    #include <iostream>

    struct PrivateTimingCPU {
        double  PCFreq;
        __int64 CounterStart;
    };

    // --- Default constructor
    TimingCPU::TimingCPU() { privateTimingCPU = new PrivateTimingCPU; (*privateTimingCPU).PCFreq = 0.0; (*privateTimingCPU).CounterStart = 0; }

    // --- Default destructor
    TimingCPU::~TimingCPU() { }

    // --- Starts the timing
    void TimingCPU::StartCounter()
    {
        LARGE_INTEGER li;
        if(!QueryPerformanceFrequency(&li)) std::cout << "QueryPerformanceFrequency failed!\n";

        (*privateTimingCPU).PCFreq = double(li.QuadPart)/1000.0;

        QueryPerformanceCounter(&li);
        (*privateTimingCPU).CounterStart = li.QuadPart;
    }

    // --- Gets the timing counter in ms
    double TimingCPU::GetCounter()
    {
        LARGE_INTEGER li;
        QueryPerformanceCounter(&li);
        return double(li.QuadPart-(*privateTimingCPU).CounterStart)/(*privateTimingCPU).PCFreq;
    }
#endif

TimingCPU.h

// 1 micro-second accuracy
// Returns the time in seconds

#ifndef __TIMINGCPU_H__
#define __TIMINGCPU_H__

#ifdef __linux__

    class TimingCPU {

        private:
            long cur_time_;

        public:

            TimingCPU();

            ~TimingCPU();

            void StartCounter();

            double GetCounter();
    };

#elif _WIN32 || _WIN64

    struct PrivateTimingCPU;

    class TimingCPU
    {
        private:
            PrivateTimingCPU *privateTimingCPU;

        public:

            TimingCPU();

            ~TimingCPU();

            void StartCounter();

            double GetCounter();

    }; // TimingCPU class

#endif

#endif

TimingGPU.cu

/**************/
/* TIMING GPU */
/**************/

#include "TimingGPU.cuh"

#include <cuda.h>
#include <cuda_runtime.h>

struct PrivateTimingGPU {
    cudaEvent_t     start;
    cudaEvent_t     stop;
};

// default constructor
TimingGPU::TimingGPU() { privateTimingGPU = new PrivateTimingGPU; }

// default destructor
TimingGPU::~TimingGPU() { }

void TimingGPU::StartCounter()
{
    cudaEventCreate(&((*privateTimingGPU).start));
    cudaEventCreate(&((*privateTimingGPU).stop));
    cudaEventRecord((*privateTimingGPU).start,0);
}

void TimingGPU::StartCounterFlags()
{
    int eventflags = cudaEventBlockingSync;

    cudaEventCreateWithFlags(&((*privateTimingGPU).start),eventflags);
    cudaEventCreateWithFlags(&((*privateTimingGPU).stop),eventflags);
    cudaEventRecord((*privateTimingGPU).start,0);
}

// Gets the counter in ms
float TimingGPU::GetCounter()
{
    float   time;
    cudaEventRecord((*privateTimingGPU).stop, 0);
    cudaEventSynchronize((*privateTimingGPU).stop);
    cudaEventElapsedTime(&time,(*privateTimingGPU).start,(*privateTimingGPU).stop);
    return time;
}

TimingGPU.cuh

#ifndef __TIMING_CUH__
#define __TIMING_CUH__

/**************/
/* TIMING GPU */
/**************/

// Events are a part of CUDA API and provide a system independent way to measure execution times on CUDA devices with approximately 0.5
// microsecond precision.

struct PrivateTimingGPU;

class TimingGPU
{
    private:
        PrivateTimingGPU *privateTimingGPU;

    public:

        TimingGPU();

        ~TimingGPU();

        void StartCounter();
        void StartCounterFlags();

        float GetCounter();

}; // TimingCPU class

#endif

  • ¡Funciona genial! Tuve que incluir también #include “TimingCPU.cpp” y #include “TimingGPU.cu” además de los dos incluidos mencionados anteriormente.

    – Gigi

    16 mayo 2020 a las 23:17


Hay un fuera de caja GpuTemporizador estructura de uso:

#ifndef __GPU_TIMER_H__
#define __GPU_TIMER_H__

struct GpuTimer
{
      cudaEvent_t start;
      cudaEvent_t stop;

      GpuTimer()
      {
            cudaEventCreate(&start);
            cudaEventCreate(&stop);
      }

      ~GpuTimer()
      {
            cudaEventDestroy(start);
            cudaEventDestroy(stop);
      }

      void Start()
      {
            cudaEventRecord(start, 0);
      }

      void Stop()
      {
            cudaEventRecord(stop, 0);
      }

      float Elapsed()
      {
            float elapsed;
            cudaEventSynchronize(stop);
            cudaEventElapsedTime(&elapsed, start, stop);
            return elapsed;
      }
};

#endif  /* __GPU_TIMER_H__ */

Si desea medir el tiempo de GPU, tiene que usar eventos. Hay un gran hilo de discusión sobre lo que se debe y lo que no se debe hacer al programar su solicitud en el foros de nvidia aquí.

Puede usar el generador de perfiles de visula computacional que será excelente para su propósito. mide el tiempo de cada función cuda y te dice cuántas veces la llamaste.

  • Gracias, pero necesito hacer estas medidas programáticamente.

    – Tudor

    24/10/2011 a las 17:13

  • @Programmer: El generador de perfiles también serializa completamente la API y agrega latencia porque requiere transferencias adicionales de dispositivos host para recopilar la salida del contador de perfiles. Es útil para muchas cosas, pero los tiempos de ejecución precisos no son una de ellas.

    – garras

    24/10/2011 a las 18:14

  • @talonmies: ¿Qué quiere decir con que el generador de perfiles serializa completamente la API? –

    – Programador

    25 de octubre de 2011 a las 9:00

  • @Programmer: la API de CUDA es naturalmente asíncrona (lanzamientos de kernel, transmisiones, ciertas clases de transferencias de memoria). Cuando ejecuta programas en el generador de perfiles, todos se vuelven seriales. Si tiene un código que se superpone a la copia de la memoria con la ejecución del kernel, serán seriales cuando se perfilen. En Fermi, la ejecución múltiple y simultánea del kernel también está deshabilitada durante la generación de perfiles.

    – garras

    25 de octubre de 2011 a las 12:04

  • Creo que aún puede hacerlo mediante programación si usa el perfilador de línea de comando (no el perfilador visual). Pero como dijo talonmies, serializa las llamadas a la API. Entonces, lo que obtienes es ejecutar todas las llamadas a la API de manera bloqueada. Y también hay una pequeña sobrecarga adicional para leer los contadores.

    – Zk1001

    4 de noviembre de 2011 a las 5:51

  • Gracias, pero necesito hacer estas medidas programáticamente.

    – Tudor

    24/10/2011 a las 17:13

  • @Programmer: El generador de perfiles también serializa completamente la API y agrega latencia porque requiere transferencias adicionales de dispositivos host para recopilar la salida del contador de perfiles. Es útil para muchas cosas, pero los tiempos de ejecución precisos no son una de ellas.

    – garras

    24/10/2011 a las 18:14

  • @talonmies: ¿Qué quiere decir con que el generador de perfiles serializa completamente la API? –

    – Programador

    25 de octubre de 2011 a las 9:00

  • @Programmer: la API de CUDA es naturalmente asíncrona (lanzamientos de kernel, transmisiones, ciertas clases de transferencias de memoria). Cuando ejecuta programas en el generador de perfiles, todos se vuelven seriales. Si tiene un código que se superpone a la copia de la memoria con la ejecución del kernel, serán seriales cuando se perfilen. En Fermi, la ejecución múltiple y simultánea del kernel también está deshabilitada durante la generación de perfiles.

    – garras

    25 de octubre de 2011 a las 12:04

  • Creo que aún puede hacerlo mediante programación si usa el perfilador de línea de comando (no el perfilador visual). Pero como dijo talonmies, serializa las llamadas a la API. Entonces, lo que obtienes es ejecutar todas las llamadas a la API de manera bloqueada. Y también hay una pequeña sobrecarga adicional para leer los contadores.

    – Zk1001

    4 de noviembre de 2011 a las 5:51

¿Ha sido útil esta solución?

Esta web utiliza cookies propias y de terceros para su correcto funcionamiento y para fines analíticos y para mostrarte publicidad relacionada con sus preferencias en base a un perfil elaborado a partir de tus hábitos de navegación. Al hacer clic en el botón Aceptar, acepta el uso de estas tecnologías y el procesamiento de tus datos para estos propósitos. Configurar y más información
Privacidad