Блог

Главная - Вопросы по программированию - Повышение производительности ядра за счет увеличения занятости?

Повышение производительности ядра за счет увеличения занятости?

#cuda

#cuda

Вопрос:

Вот вывод Compute Visual Profiler для моего ядра на GT 440:

  • Сведения о ядре: размер сетки: [100 1 1], размер блока: [256 1 1]
  • Соотношение регистров: 0,84375 ( 27648 / 32768) [35 регистров на поток]
  • Коэффициент разделяемой памяти: 0,336914 ( 16560/49 152) [5520 байт на блок]
  • Активных блоков на SM: 3 (Максимальное количество активных блоков на SM: 8)
  • Активных потоков на SM: 768 (Максимальное количество активных потоков на SM: 1536)
  • Потенциальная занятость: 0,5 ( 24/48)
  • Фактор, ограничивающий занятость: регистры

Пожалуйста, обратите внимание на маркеры, выделенные жирным шрифтом. Время выполнения ядра составляет 121195 us .

Я уменьшил количество регистров в потоке, переместив некоторые локальные переменные в общую память. Выходные данные Compute Visual Profiler стали:

  • Сведения о ядре: размер сетки: [100 1 1], размер блока: [256 1 1]
  • Соотношение регистров: 1 ( 32768 / 32768) [30 регистров на поток]
  • Коэффициент разделяемой памяти: 0,451823 ( 22208 / 49152) [5552 байта на блок]
  • Активных блоков на SM: 4 (Максимальное количество активных блоков на SM: 8)
  • Активных потоков на SM: 1024 (Максимальное количество активных потоков на SM: 1536)
  • Потенциальная занятость: 0,666667 ( 32/48 )
  • Фактор, ограничивающий занятость: регистры

Следовательно, теперь 4 блоки выполняются одновременно на одном SM по сравнению 3 с блоками в предыдущей версии. Тем не менее, время выполнения 115756 us почти такое же! Почему? Разве блоки не выполняются полностью независимо на разных ядрах CUDA?

Ответ №1:

Вы неявно предполагаете, что более высокая занятость автоматически приводит к повышению производительности. Чаще всего это не так.

Архитектуре NVIDIA требуется определенное количество активных деформаций на МП, чтобы скрыть задержку конвейера команд графического процессора. На вашей плате на базе Fermi это требование означает минимальную загрузку около 30%. Стремление к увеличению занятости, превышающей этот минимум, не обязательно приведет к увеличению пропускной способности, поскольку узкое место задержки может переместиться в другую часть графического процессора. Ваш графический процессор начального уровня не обладает большой пропускной способностью памяти, и вполне возможно, что 3 блока на МП достаточно, чтобы ограничить пропускную способность памяти кода, и в этом случае увеличение количества блоков не окажет никакого влияния на производительность (оно может даже снизиться из-за увеличения конкуренции с контроллером памятии промахи в кэше). Кроме того, вы сказали, что передали переменные в разделяемую память, чтобы уменьшить размер регистра ядра. В Fermi общая память имеет пропускную способность всего около 1000 Гбит / с по сравнению с примерно 8000 Гбит / с для регистров (см. Ссылку ниже для результатов микропомаркирования, которые демонстрируют это). Итак, вы переместили переменные в более медленную память, что также может негативно сказаться на производительности, сводя на нет любые преимущества, которые дает высокая занятость.

Если вы еще не видели его, я настоятельно рекомендую презентацию Василия Волкова из GTC 2010 «Лучшая производительность при меньшей заполняемости» (pdf). Здесь показано, как использование параллелизма на уровне команд может увеличить пропускную способность графического процессора до очень высоких уровней при очень, очень низких уровнях загрузки.

Комментарии:

1. Хороший ответ. Занятость является серьезной проблемой только для сокрытия глобальной задержки доступа к памяти; для потоков, связанных с вычислениями, нескольких активных потоков на SP должно быть достаточно. Это тоже ваше понимание?

2. Я действительно так не думаю, Патрик. Это не относится ко всем типам ядер. Для ядер, связанных с вычислениями, более высокая занятость может все же увеличить производительность. Не так просто определить, сколько активных искажений необходимо для сокрытия арифметической задержки. Это зависит от типов операций и того, как они чередуются друг с другом.

Ответ №2:

talonmies уже ответил на ваш вопрос, поэтому я просто хочу поделиться кодом, вдохновленным первой частью презентации В. Волкова, упомянутой в ответе выше.

Это код:

 #include<stdio.h>

#define N_ITERATIONS 8192

//#define DEBUG

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %dn", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/********************************************************/
/* KERNEL0 - NO INSTRUCTION LEVEL PARALLELISM (ILP = 0) */
/********************************************************/
__global__ void kernel0(int *d_a, int *d_b, int *d_c, unsigned int N) {

    const int tid = threadIdx.x   blockIdx.x * blockDim.x ;

    if (tid < N) {

        int a = d_a[tid];
        int b = d_b[tid];
        int c = d_c[tid];

        for(unsigned int i = 0; i < N_ITERATIONS; i  ) {
            a = a * b   c;
        }

        d_a[tid] = a;
    }

}

/*****************************************************/
/* KERNEL1 - INSTRUCTION LEVEL PARALLELISM (ILP = 2) */
/*****************************************************/
__global__ void kernel1(int *d_a, int *d_b, int *d_c, unsigned int N) {

    const int tid = threadIdx.x   blockIdx.x * blockDim.x;

    if (tid < N/2) {

        int a1 = d_a[tid];
        int b1 = d_b[tid];
        int c1 = d_c[tid];

        int a2 = d_a[tid N/2];
        int b2 = d_b[tid N/2];
        int c2 = d_c[tid N/2];

        for(unsigned int i = 0; i < N_ITERATIONS; i  ) {
            a1 = a1 * b1   c1;
            a2 = a2 * b2   c2;
        }

        d_a[tid]        = a1;
        d_a[tid N/2]    = a2;
    }

}

/*****************************************************/
/* KERNEL2 - INSTRUCTION LEVEL PARALLELISM (ILP = 4) */
/*****************************************************/
__global__ void kernel2(int *d_a, int *d_b, int *d_c, unsigned int N) {

    const int tid = threadIdx.x   blockIdx.x * blockDim.x;

    if (tid < N/4) {

        int a1 = d_a[tid];
        int b1 = d_b[tid];
        int c1 = d_c[tid];

        int a2 = d_a[tid N/4];
        int b2 = d_b[tid N/4];
        int c2 = d_c[tid N/4];

        int a3 = d_a[tid N/2];
        int b3 = d_b[tid N/2];
        int c3 = d_c[tid N/2];

        int a4 = d_a[tid 3*N/4];
        int b4 = d_b[tid 3*N/4];
        int c4 = d_c[tid 3*N/4];

        for(unsigned int i = 0; i < N_ITERATIONS; i  ) {
            a1 = a1 * b1   c1;
            a2 = a2 * b2   c2;
            a3 = a3 * b3   c3;
            a4 = a4 * b4   c4;
        }

        d_a[tid]        = a1;
        d_a[tid N/4]    = a2;
        d_a[tid N/2]    = a3;
        d_a[tid 3*N/4]  = a4;
    }

}

/********/
/* MAIN */
/********/
void main() {

    const int N = 1024;

    int *h_a                = (int*)malloc(N*sizeof(int));
    int *h_a_result_host    = (int*)malloc(N*sizeof(int));
    int *h_a_result_device  = (int*)malloc(N*sizeof(int));
    int *h_b                = (int*)malloc(N*sizeof(int));
    int *h_c                = (int*)malloc(N*sizeof(int));

    for (int i=0; i<N; i  ) {
        h_a[i] = 2;
        h_b[i] = 1;
        h_c[i] = 2;
        h_a_result_host[i] = h_a[i];
        for(unsigned int k = 0; k < N_ITERATIONS; k  ) {
            h_a_result_host[i] = h_a_result_host[i] * h_b[i]   h_c[i];
        }
    }

    int *d_a; gpuErrchk(cudaMalloc((void**)amp;d_a, N*sizeof(int)));
    int *d_b; gpuErrchk(cudaMalloc((void**)amp;d_b, N*sizeof(int)));
    int *d_c; gpuErrchk(cudaMalloc((void**)amp;d_c, N*sizeof(int)));

    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_b, h_b, N*sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_c, h_c, N*sizeof(int), cudaMemcpyHostToDevice));

    // --- Creating events for timing
    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(amp;start);
    cudaEventCreate(amp;stop);

    /***********/
    /* KERNEL0 */
    /***********/
    cudaEventRecord(start, 0);
    kernel0<<<1, N>>>(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(amp;time, start, stop);
    printf("GFlops = %fn", (1.e-6)*(float)(N*N_ITERATIONS)/time);
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost));
    for (int i=0; i<N; i  ) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %in", i, h_a_result_host[i], h_a_result_device[i]); return; }

    /***********/
    /* KERNEL1 */
    /***********/
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice));
    cudaEventRecord(start, 0);
    kernel1<<<1, N/2>>>(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(amp;time, start, stop);
    printf("GFlops = %fn", (1.e-6)*(float)(N*N_ITERATIONS)/time);
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost));
    for (int i=0; i<N; i  ) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %in", i, h_a_result_host[i], h_a_result_device[i]); return; }

    /***********/
    /* KERNEL2 */
    /***********/
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice));
    cudaEventRecord(start, 0);
    kernel2<<<1, N/4>>>(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(amp;time, start, stop);
    printf("GFlops = %fn", (1.e-6)*(float)(N*N_ITERATIONS)/time);
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost));
    for (int i=0; i<N; i  ) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %in", i, h_a_result_host[i], h_a_result_device[i]); return; }

    cudaDeviceReset();

}

На моей GeForce GT540M результат

 kernel0   GFlops = 21.069281    Occupancy = 66%
kernel1   GFlops = 21.183354    Occupancy = 33%
kernel2   GFlops = 21.224517    Occupancy = 16.7%

это означает, что ядра с меньшей занятостью все еще могут демонстрировать высокую производительность, если используется параллелизм на уровне команд (ILP).