Medindo e comparando precisamente o tempo de execução

[Alexsander Petry]

Para medir e comparar os tempos de execução na CPU e na GPU, precisamos medir o tempo da MESMA forma, senão a comparação não é válida.
Essa medição não pode ser baseada no tempo global do sistema, pois o sistema operacional pode escalonar outras tarefas durante a execução do seu kernel/programa, e medindo o tempo global, o tempo de execução destas outras tarefas será contabilizado no tempo de execução do seu kernel/programa, certo?
Como vocês costumam resolver isto?
Como vocês costumam medir o tempo na CPU? E na GPU?

Grato.
--
Alexsander Petry

[victor...@gmail.com]

oi Alexsander,

sua pergunta faz muito sentido e já tive bastantes problemas com isso na hora hora de comparar o desempenho de algoritmos na gpu e cpu.

pelo que eu entendi, a pergunta é referente à interação entre seu programa em cuda e a cpu, e seus efeitos no desempenho. pois bem, em aplicações que rodam na cpu (multithread, serial), scheduling é um problema, pois nunca se sabe o tempo de execução real no processador.

nesse caso, existem ferramentas (da intel por exemplo) que se propõem a fazer a contagem só do tempo que um processo ficou de fato no processador, mas você pode ver que não é tarefa simples e quase sempre é feita no hardware mesmo. o que você pode fazer nesse caso é fazer várias execuções e tirar o _mínimo_ (não a média) de todas elas, pois esta seria a que menos teve interferência de outras aplicações, mas não vá nessa se puder, pois mesmo assim os resultados podem ser errados. o melhor é fazer benchmarks com o pc o mais livre possível.

na gpu, temos uma situação estranha, a maior parte das gpus nvidia só executa kernels de um único processo por vez (podemos intercalar kernels usando filas e tudo o mais, mas continuam sendo de um único processo) e é possível fazer a medição de tempo no seu programa main com gettimeofday ou cuda events. mas, existem algumas gpus já no mercado (http://www.nvidia.com/object/sli_multi_os.html) que permitem a virtualização do device, i.e. várias  compartilhando uma gpu. suponho que no futuro isso vai se tornar comum e aí vai ser o caos para fazer as medições :)

a observação do andré quanto ao uso da placa para o monitor é importante também, não sei os efeitos disso no desempenho, em geral eu faço os testes em modo texto no linux (e só assim dá para sair do limite máximo de 5seg de execução de um kernel com uma gpu apenas).

ou seja, dá sim para contar o tempo usando cuda events, mas não sei como vai ser no futuro.

2011/2/24 André Körbes <andre....@gmail.com>

Olá Alexsander,

Em geral este tipo de consideração não é muito vista por aí.

Não sei se algum sistema operacional permite que vc defina que ele deve executar apenas uma tarefa, então vc pode contar que sempre vai ter o efeito do escalonamento. O que eu vejo que é feito a respeito disso é executar o mesmo programa diversas vezes, para se obter uma média que de certa forma anule o efeito do escalonamento de tarefas. Logicamente, não rodar mais nada além do SO no PC ao mesmo tempo é uma boa ideia.

Sobre o tempo na GPU, há mais de um problema além do compartilhamento. Você tem que considerar a transferência de memória se quiser uma comparação justa, pois na CPU os dados já estão lá. Além disso, se vc usa a mesma GPU para o vídeo também está ocupando tempo de processamento.

Acho que aqui vale a mesma coisa, rodar o mesmo programa várias vezes para obter uma média consistente.

Uma questão que também vejo é que, em casos reais, seu algoritmo vai rodar sendo escalonado por um sistema operacional, então não sei se faz muito sentido rodar ele sozinho, se é que isso é possível. Acho essa discussão muito interessante, e gostaria de ouvir a opinião dos outros colegas da lista.

Abraços,

André

2011/2/24 Alexsander Petry <alexsan...@gmail.com>

--
Você está recebendo esta mensagem porque se inscreveu no grupo "GPU Brasil" dos Grupos do Google.
Para postar neste grupo, envie um e-mail para gpub...@googlegroups.com.
Para cancelar a inscrição nesse grupo, envie um e-mail para gpubrasil+...@googlegroups.com.
Para obter mais opções, visite esse grupo em http://groups.google.com/group/gpubrasil?hl=pt-BR.

--
Você está recebendo esta mensagem porque se inscreveu no grupo "GPU Brasil" dos Grupos do Google.
Para postar neste grupo, envie um e-mail para gpub...@googlegroups.com.
Para cancelar a inscrição nesse grupo, envie um e-mail para gpubrasil+...@googlegroups.com.
Para obter mais opções, visite esse grupo em http://groups.google.com/group/gpubrasil?hl=pt-BR.

[André Körbes]

Olá Alexsander,

Em geral este tipo de consideração não é muito vista por aí.

Não sei se algum sistema operacional permite que vc defina que ele deve executar apenas uma tarefa, então vc pode contar que sempre vai ter o efeito do escalonamento. O que eu vejo que é feito a respeito disso é executar o mesmo programa diversas vezes, para se obter uma média que de certa forma anule o efeito do escalonamento de tarefas. Logicamente, não rodar mais nada além do SO no PC ao mesmo tempo é uma boa ideia.

Sobre o tempo na GPU, há mais de um problema além do compartilhamento. Você tem que considerar a transferência de memória se quiser uma comparação justa, pois na CPU os dados já estão lá. Além disso, se vc usa a mesma GPU para o vídeo também está ocupando tempo de processamento.

Acho que aqui vale a mesma coisa, rodar o mesmo programa várias vezes para obter uma média consistente.

Uma questão que também vejo é que, em casos reais, seu algoritmo vai rodar sendo escalonado por um sistema operacional, então não sei se faz muito sentido rodar ele sozinho, se é que isso é possível. Acho essa discussão muito interessante, e gostaria de ouvir a opinião dos outros colegas da lista.

Abraços,

André

2011/2/24 Alexsander Petry <alexsan...@gmail.com>

Para medir e comparar os tempos de execução na CPU e na GPU, precisamos medir o tempo da MESMA forma, senão a comparação não é válida.

--

[Alexsander Petry]

Obrigado por responderem, também acho esta questão muito importante, e seria interessante se ao final da discussão tivessemos publicado, aqui na lista, um padrão confiável e preciso (ou o mais próximo disso) de medir os tempos de execução na CPU e GPU.

Minha preocupação não é encontrar um Sistema Operacional monotarefa, mas sim uma maneira de medir o tempo que meu algoritmo realmente passou na CPU, no linux, é possível fazer esta medição, usando as funções times, clock, clock_gettime, etc. Sua sugestão de tirar a média dos tempos em execução é válida, já vi dissertações usando este método, mas não sei dizer se é a mais correta.

Acredito que uma medição válida e confiável tenha que ser feita da mesma forma em algoritmos processados na CPU e na GPU+CPU (pois a parte de comunicação CPU <--> GPU também faz parte do tempo de processamento do algoritmo). Então por exemplo: Em um algoritmo de compressão de video, na CPU o tempo deve ser medido apenas o tempo que ele realmente ficou no processador, na GPU deve ser contado o tempo para preparar os dados para serem enviados à GPU MAIS o tempo de transmissão da CPU para a GPU MAIS o tempo de execução na própria MAIS o tempo de transmissão de volta da GPU para a CPU MAIS o tempo de preparação dos dados para a aplicação, assim teremos isolado o tempo de execução e o de preparação e transmissão dos dados. Acredito que é importante medir o tempo de transmissão e preparação dos dados, pois na CPU isto não é preciso, então ele deve ser contabilizado no tempo total de execução do algoritmo.

Concordo que para esta medição, a GPU deve estar totalmente disponível para o processamento do algoritmo, ou seja, sem estar no modo gráfico, e na execução na CPU, o sistema operacional também deve ser executado em modo texto e deve-se deixá-lo o mais livre possível, ou seja, sem "nada" instalado e sem "nada" em execução.

Sendo assim, e considerando que faremos a medição em linux, então:

 - Na CPU podemos medir apenas o tempo que ele ficou no processador, usando as funções que citei acima. Para medir o tempo de execução.

 - Na GPU, para medir o tempo que o algoritmo ficou na GPU mais os tempos das transmissões, podemos usar o tempo global do sistema. Para medir o tempo de preparação dos dados, podemos medir e tirar diferenças dos tempos de CPU.

O que vocês acham?

Mas ainda não sei se este método é totalmente confiável, pois não sei em que estado a CPU fica enquanto o algoritmo está sendo processado na GPU e durante a transmissão. O processo bloqueia a CPU até a resposta da GPU? O processo fica "dormindo" até a resposta da GPU? Ou o que? Alguém sabe?

Alguém se aventura a postar um primeiro código exemplo?

Grato.

Alexsander Petry

[victor...@gmail.com]

quando você executa um kernel, o controle volta imediatamente para a cpu (i.e. assincrono), mas quando você faz alguma operação como cudaMemcpy, cudaThreadsSynchronize, etc. há uma sincronização entre os dois.

inclusive isso é usado para intercalar execuções na cpu e na gpu.

2011/2/24 Alexsander Petry <alexsan...@gmail.com>

[Thársis Tuani Pinto Souza]

Exemplo básico de código que mede "desempenho" de execução de kernel:

http://developer.download.nvidia.com/compute/cuda/2_3/sdk/projects/clock.tar.gz

O tempo de execução é medido por bloco.

--
Thársis T. P. Souza
Engenheiro de Computação - UNICAMP
Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/2/25 victor...@gmail.com <victor...@gmail.com>

[Thársis Tuani Pinto Souza]

CUDA profiler:

http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/visual_profiler_cuda/CUDA_Profiler_3.0.txt

--
Thársis T. P. Souza
Engenheiro de Computação - UNICAMP
Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/2/25 Thársis Tuani Pinto Souza <souza....@gmail.com>

[Eduardo - Yahoo]

Gostei do rumo que a discussão tomou mas continuo com a mesma dúvida.
Como mensurar o tempo de execução de cada thread de um kernel?

Atenciosamente,

Eduardo Camargo
Analista de Sistemas

http://www.lncc.br/mcs/
HeMoLab - http://hemolab.lncc.br/equipe/desenv/eduardo.htm
Laboratório Nacional de Computação Científica - LNCC

[Thársis Tuani Pinto Souza]

Em teoria, teria que ser com a CUDA event API. Abaixo um exemplo referente a kernel, acredito ser análogo a thread.

// Prepare
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);


// Start record
cudaEventRecord(start, 0);
// Do something on GPU
MyKernel<<<dimGrid, dimBlock>>>(input_data, output_data);
// Stop event
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);


float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop); // that's our time!
// Clean up:
cudaEventDestroy(start);
cudaEventDestroy(stop);

--
Thársis T. P. Souza
Engenheiro de Computação - UNICAMP
Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/2/26 Eduardo - Yahoo <camar...@yahoo.com.br>

[Vinicius Rocha]

Nesse caso o que está sendo medido é o tempo de execução do kernel, e não de cada thread individualmente.

2011/2/26 Thársis Tuani Pinto Souza <souza....@gmail.com>

--
Vinicius Rocha

MCPD | Microsoft Student Partner

System Engineer @ c.e.s.a.r 

[Dinart Duarte Braga]

Verdade Vinicius,

Talvez se ele chamar um kernel só com uma thread, dê para ter uma ideia do tempo de execução de uma única thread.

Em tese, se ele chamar esse kernel com um número de threads menor que o número máximo de threads que a GPU dele pode executar simultaneamente, ele terá o tempo de execução de uma thread.

É claro que essa consideração deixa de ser válida se as threads não forem homogêneas, ou seja, não executarem exatamente as mesmas instruções, devido a presença de um ramo condicional.

Dinart Duarte Braga

[Alexsander Petry]

Tharsis, obrigado pelo código exemplo, não cheguei a testá-lo, mas parece que funciona.
Seria legal fazer uma comparação do tempo medido por este código e do diff do tempo global do sistema (por exemplo a função gettimeofday), para ver se há grande diferença.
Não posso fazer agora pois estou reinstalando tudo aqui, se alguém conseguir fazer este teste e postar aqui, seria muito interessante.

Medir o tempo de execução total de um kernel, no meu caso, é mais útil do que medir o tempo de cada thread, mas as duas discuções são interessantes e úteis.

[Eduardo - Yahoo]

Tharsis, vendo seu código imagino de vc deve ter imaginado que colocando as chamadas ao CUDA Event API dentro no Kernel (no início e no final) estaria medindo o tempo de cada thread. Meu raciocínio tb me levou à isto mas o que acontece é que a execução não passa das chamadas iniciais.
Era isso mesmo que sugeriu? Alguma idéia de como fazer o que quero?


Saudações à todos,
Eduardo

2011/2/26 Alexsander Petry <alexsan...@gmail.com>

[Thársis Tuani Pinto Souza]

Para medir o tempo para cada thread (parametrizada), acredito que seria necessário parametrizar analogamente a medição.

[Eduardo - Yahoo]

Desculpe a ignorância mas poderia me explicar melhor?

[Thársis Tuani Pinto Souza]

Uma sugestão seria fazer um vetor de cudaEvent parametrizado conforme sua thread.
Ex.:

cudaEventRecord(start[blockIdx.x * blockDim.x + threadIdx.x], 0);



...
código da thread[blockIdx.x * blockDim.x + threadIdx.x]


...

cudaEventRecord(stop[blockIdx.x * blockDim.x + threadIdx.x], 0);


cudaEventSynchronize(stop[blockIdx.x * blockDim.x + threadIdx.x]);

Mas realmente é complicado medir o tempo da thread. Se alguém tiver outra sugestão manda aí!

[Eduardo - Yahoo]

Gostei da sugestão, irei testar e posto o resultado.


Obrigado.

[Thársis Tuani Pinto Souza]

Valeu!

Essa discussão é bem pertinente e recorrente para quem necessita ser rigoroso com desempenho.

[Eduardo - Yahoo]

Amigos,

Tentei fazer a medição do tempo de execução de cada thread utilizando as idéias que me deram e não consegui. Estou começando a acreditar que a CUDA não permite este tipo de medição, ao menos usando CUDAEvent. A mensagem retornada (esta abaixo) no momento da compilação sugere que os métodos cudaEventCreate, cudaEventRecord, cudaEventSynchronize, cudaEventElapsedTime, cudaEventDestroy só podem ser executados no host. Esta mensagem é informada referenciando cada linha onde os métodos aparecem.


calling a host function from a __device__/__global__ function is not allowed


Muito obrigado à todos que estão me ajudando com suas sugestões e idéias!!!

Abs
Eduardo

[Edward Boszczowski]

Uma sugestão rápida.
Alternativamente já tentou postar este assunto no fórum CUDA (http://forums.nvidia.com/index.php?s=bbf4f13859f5936702719bb0fe666e8c&showforum=64)?

Creio que lá eles também podem dar alguma ajuda.

http://about.me/edward_boszczowski

[Alexsander Petry]

Fiz um código para medir o tempo de processamento de um certo algoritmo na cpu.
Basicamente é pegar o tempo antes no inicio e no fim do processamento, e tirar a diferença.
O código calcula os tempos de duas maneiras diferentes, uma usando o tempo global do sistema (gettimeofday) e outra usando o tempo de cpu do processo (clock_gettime).
Fazendo as medições desta maneira, cheguei a um resultado válido e confiável.
Se alguém tiver alguma sugestão de melhoria, será bem vinda.

Se alguém quiser fazer algo parecido para medir o tempo de execução de um algoritmo na GPU, seria muito interessante.

Este código foi feito para linux, não cheguei a testá-lo em windows.
O arquivo se chama tempocpu.c.
Para compilá-lo, execute: gcc -Wall -lrt tempocpu.c -o tempocpu

Segue o código...

// Direitos: Alexsander Petry - 03/02/2011

#include <stdio.h>
#include <string.h>

#include <sys/time.h>   // gettimeofday  : dif global
#include <time.h>       // clock_gettime : dif tempo de cpu

#include <unistd.h>     // usleep

long fatorial(int num)
{
   int i=0;
   long fat=1;
   for(i = 1; i <= num; i++) {
      fat *= i;
   }
   return fat;
}

void proc_math()
{
   int i=0, ii=0;
   for(i = 1; i <= 10; i++) {
      for(ii = 1; ii <= 1000; ii++) {
         if(fatorial(100) != fatorial(100)) {
            printf("ERRO: fatorial != fatorial ???? \n");
         }
      }
   }
}

int main()
{
   struct timeval tvi, tve;         // gettimeofday : dif global
   long tv_dif_us=0, tv_dif_ms=0;   // gettimeofday : dif global

   struct timespec tsi, tse;        // clock_gettime : dif tempo de cpu
   long ts_dif_us=0, ts_dif_ms=0;   // clock_gettime : dif tempo de cpu
  
   if(gettimeofday(&tvi, NULL) != 0) {
      printf("ERRO no gettimeofday \n");
      return 1;
   }
   if(clock_gettime(CLOCK_PROCESS_CPUTIME_ID /*CLOCK_REALTIME*/, &tsi) != 0) {
      printf("ERRO no clock_gettime \n");
      return 1;
   }
  
   proc_math(); // rotina para ocupar a cpu
   // usleep(1 * 1000/*ms*/ * 1000/*seg*/);
  
   if(gettimeofday(&tve, NULL) != 0) {
      printf("ERRO no gettimeofday \n");
      return 1;
   }
   if(clock_gettime(CLOCK_PROCESS_CPUTIME_ID /*CLOCK_REALTIME*/, &tse) != 0) {
      printf("ERRO no clock_gettime \n");
      return 1;
   }
  
   tv_dif_us = (tve.tv_sec - tvi.tv_sec) * 1000 * 1000 + (tve.tv_usec - tvi.tv_usec);
   tv_dif_ms = tv_dif_us / 1000;
   printf("---- Dif global ---- \n");
   printf("Ini: %10lds:%6ldus \n", tvi.tv_sec, tvi.tv_usec);
   printf("End: %10lds:%6ldus \n", tve.tv_sec, tve.tv_usec);
   printf("Dif: %10lds:%6ldus \n", tv_dif_ms / 1000, tv_dif_us - (tv_dif_ms / 1000) * 1000);
   printf("Dif total: %ldus ou %ldms \n", tv_dif_us, tv_dif_ms);
  
   ts_dif_us = (tse.tv_sec - tsi.tv_sec) * 1000 * 1000 + (tse.tv_nsec - tsi.tv_nsec) / 1000; // nsec / 1000 para converter para usec
   ts_dif_ms = ts_dif_us / 1000;
   printf("---- Dif tempo cpu ---- \n");
   printf("Ini: %10lds:%6ldus \n", tsi.tv_sec, tsi.tv_nsec / 1000);
   printf("End: %10lds:%6ldus \n", tse.tv_sec, tse.tv_nsec / 1000);
   printf("Dif: %10lds:%6ldus \n", ts_dif_ms / 1000, ts_dif_us - (ts_dif_ms / 1000) * 1000);
   printf("Dif total: %ldus ou %ldms \n", ts_dif_us, ts_dif_ms);

   return 0;

[Thársis Tuani Pinto Souza]

Valeu pela colaboração Alexsander!

Em anexo, envio um .c que mede tempo de cpu utilizado por ticks de relogio transcorridos e um .cu com uma medição simples de tempo em GPU utilizando cudaEvent.

Abraços.

--
Thársis T. P. Souza
Engenheiro de Computação - UNICAMP
Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/3/3 Alexsander Petry <alexsan...@gmail.com>

[Alexsander Petry]

Opa! Valeu Tharsis!
Durante meus testes para chegar ao código que postei, também usei a função clock, mas acabei optando pela clock_gettime, pois já entrega o tempo em nseg, e a conversão para useg e mseg é mais simples.
Digo mais simples, porque pelo que entendi do seu cpuTime.c, ele mede o tempo em segundos, certo? Então para converter para useg ou mseg, ao invés de dividir o número de ticks por CLOCKS_PER_SEC, deveriamos dividir por
(CLOCKS_PER_SEC / (1000 * 1000)), para obter o resultado em useg, certo?. Isto pode não parecer muito claro para a maioria.

Vou tentar usar a ideia do seu cudaEvent.cu para medir o tempo do meu algoritmo na GPU, em breve posto o resultado e comparação.

[Alexsander Petry]

Colando aqui os links dos videos do seminário apresentado pelo Thársis, que trata justamente sobre este tema e que apresenta uma conclusão sobre.
Resumindo: Prefira utilizar o tempo de processamento para medir desempenho.

http://www.youtube.com/watch?v=M0ExrGuk-Ms
http://www.youtube.com/watch?v=Zrwo6efrTsw

Estes e outros videos estão disponíveis no canal no youtube do GPUBrasil. http://www.youtube.com/user/gpubrasildotcom

--
Alexsander Petry

[Thársis T. P. Souza]

daria pra fazer um seminário SOMENTE sobre medir desempenho!
nesse seminário, apenas introduzi o assunto antes de falar de outras considerações como

1.pinned memory,
2.overlap em transferencia de dados e
3.kernels paralelos

Att.

--
Thársis T. P. Souza

Engenheiro de Computação - Unicamp

Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/5/13 Alexsander Petry <alexsan...@gmail.com>

[Alexsander Petry]

Acredito que sim!
Por exemplo, fiquei sem entender o que um cara da platéia falou sobre ele utilizar o clock da GPU para analisar se o gargalo da aplicação dele estava no acesso a memória ou no processamento propriamente dito.
Não entendi o que ele falou, mas pareceu interessante.

[Thársis T. P. Souza]

na verdade, a idéia é bem simples: diminuir o clock do processamento em x% depois medir a redução de desempenho y% e depois fazer o mesmo no acesso à memória e, com isso, concluir se o aplicativo em questão é memory-bound ou compute-bound

Att.

--
Thársis T. P. Souza
Engenheiro de Computação - Unicamp
Mestrando em Ciência da Computação - USP
Analista de Sistemas - Centro de Computação Eletrônica - USP

2011/5/14 Alexsander Petry <alexsan...@gmail.com>