dram_write_bytes результат на P100 - PullRequest
Новая
       80

dram_write_bytes результат на P100

0 голосов
/ 14 июля 2020

Я использовал nvprof для профилирования простого примера vecadd (n = 1024) на P100, но заметил, что dram_write_bytes составляет всего 256 (а не 1024 * 4, как я ожидал). Может кто-нибудь объяснить, почему это число невелико? Какие еще метрики мне нужно добавить для подсчета операций записи в глобальную память? Спасибо. Число float_count_sp верное (1024). 

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

__global__ void vecAdd(float* a, float* b, float* c, int n){
    int id = blockIdx.x*blockDim.x + threadIdx.x;
    if(id < n) c[id] = a[id] + b[id];
}

int main(int argc, char* argv[]){
    int n = 1024;
    float *h_a, *d_a;
    float *h_b, *d_b;
    float *h_c, *d_c;
    size_t bytes = n*sizeof(float);
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    int i;
    for(i = 0; i < n; i++){
        h_a[i] = sin(i)*sin(i);
        h_b[i] = cos(i)*cos(i+1);
    }
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    vecAdd <<<1, 1024>>> (d_a, d_b, d_c, n);
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    float sum = 0;
    for(i = 0; i < n; i++)
        sum += h_c[i] - h_a[i] - h_b[i];
    printf("final diff: %f\n", sum/n);
    
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

Это связано с выборкой nvprof? Один раз получаю 384 байта. Иногда даже получалось 0 байт. Странно: если я изменю n на 1024 * 1024, я получу байтов больше, чем ожидал (4688032). 4688032/1024/1024/4 = 1,11.

1 Ответ

2 голосов
/ 14 июля 2020

Есть несколько причин, по которым ваши ожидания не соблюдаются и данные меняются:

  1. Система памяти GPU используется всеми движками. Основным движком является движок графики / вычислений, но другие движки, такие как движки копирования, отображение и т. Д. c. доступ к памяти устройства и счетчики управления памятью (FB = framebuffer) не имеют метода отслеживания инициатора запроса.

  2. Внедрение NVPROF не пытается удалить всю контекстную память из кеша L2 . CudaMemcpys до запуска и код воспроизведения ядра в nvprof оставят кеш L2 в несогласованном состоянии.

  3. Начальный размер 4 КБ слишком мал для точного отслеживания. Полный набор данных может быть в L2 либо из cudaMemcpy, либо из воспроизведения. Кроме того, байты, которые вы видите, могут быть получены от других клиентов, таких как кеши констант.

Настоятельно рекомендуется масштабировать размер буфера до разумного размера. На новых графических процессорах профилировщик Nsight Compute улучшил разбивку по уровням L2 для различных клиентов, чтобы помочь обнаружить неожиданный трафик c. Кроме того, Nsight Compute replay logi c очищает кэш L2, чтобы каждый повтор имел согласованное начальное состояние.

Если у вас есть подключенный монитор, рекомендуется переместить его на другой графический процессор при просмотре DRAM. счетчики. Счетчики nvprof L2 обычно фильтруют счетчик по трафику c от SM, поэтому трафик c от копирующих механизмов, контроллера дисплея, MMU, кэшей констант и т. д. c. не будет отображаться на счетчиках L2.

...