CUDA, больше потоков для той же работы = Более длительное время работы, несмотря на лучшую загруженность, Почему? - PullRequest
Новая
       10

CUDA, больше потоков для той же работы = Более длительное время работы, несмотря на лучшую загруженность, Почему?

2 голосов
/ 15 марта 2010

Я столкнулся со странной проблемой, когда увеличение занятости за счет увеличения количества потоков снижает производительность.

Я создал следующую программу для иллюстрации проблемы: 

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil.h>

__global__ void less_threads(float * d_out) {
    int num_inliers;
    for (int j=0;j<800;++j) {
        //Do 12 computations
        num_inliers += j*(j+1);
        num_inliers += j*(j+2);
        num_inliers += j*(j+3);
        num_inliers += j*(j+4);
        num_inliers += j*(j+5);
        num_inliers += j*(j+6);
        num_inliers += j*(j+7);
        num_inliers += j*(j+8);
        num_inliers += j*(j+9);
        num_inliers += j*(j+10);
        num_inliers += j*(j+11);
        num_inliers += j*(j+12);
    }

    if (threadIdx.x == -1)
        d_out[threadIdx.x] = num_inliers;
}

__global__ void more_threads(float *d_out) {
    int num_inliers;
    for (int j=0;j<800;++j) {
        // Do 4 computations
        num_inliers += j*(j+1);
        num_inliers += j*(j+2);
        num_inliers += j*(j+3);
        num_inliers += j*(j+4);
    }
    if (threadIdx.x == -1)
        d_out[threadIdx.x] = num_inliers;
}


int main(int argc, char* argv[])
{
    float *d_out = NULL;
    cudaMalloc((void**)&d_out,sizeof(float)*25000);

    more_threads<<<780,128>>>(d_out);
    less_threads<<<780,32>>>(d_out);


    return 0;
}

И выход PTX: 

    .entry _Z12less_threadsPf (
        .param .u32 __cudaparm__Z12less_threadsPf_d_out)
    {
    .reg .u32 %r<35>;
    .reg .f32 %f<3>;
    .reg .pred %p<4>;
    .loc    17  6   0
 //   2  #include <stdlib.h>
 //   3  #include <cuda_runtime.h>
 //   4  #include <cutil.h>
 //   5  
 //   6  __global__ void less_threads(float * d_out) {
$LBB1__Z12less_threadsPf:
    mov.s32     %r1, 0;
    mov.s32     %r2, 0;
    mov.s32     %r3, 0;
    mov.s32     %r4, 0;
    mov.s32     %r5, 0;
    mov.s32     %r6, 0;
    mov.s32     %r7, 0;
    mov.s32     %r8, 0;
    mov.s32     %r9, 0;
    mov.s32     %r10, 0;
    mov.s32     %r11, 0;
    mov.s32     %r12, %r13;
    mov.s32     %r14, 0;
$Lt_0_2562:
 //<loop> Loop body line 6, nesting depth: 1, iterations: 800
    .loc    17  10  0
 //   7     int num_inliers;
 //   8     for (int j=0;j<800;++j) {
 //   9         //Do 12 computations
 //  10         num_inliers += j*(j+1);
    mul.lo.s32  %r15, %r14, %r14;
    add.s32     %r16, %r12, %r14;
    add.s32     %r12, %r15, %r16;
    .loc    17  11  0
 //  11         num_inliers += j*(j+2);
    add.s32     %r17, %r15, %r12;
    add.s32     %r12, %r1, %r17;
    .loc    17  12  0
 //  12         num_inliers += j*(j+3);
    add.s32     %r18, %r15, %r12;
    add.s32     %r12, %r2, %r18;
    .loc    17  13  0
 //  13         num_inliers += j*(j+4);
    add.s32     %r19, %r15, %r12;
    add.s32     %r12, %r3, %r19;
    .loc    17  14  0
 //  14         num_inliers += j*(j+5);
    add.s32     %r20, %r15, %r12;
    add.s32     %r12, %r4, %r20;
    .loc    17  15  0
 //  15         num_inliers += j*(j+6);
    add.s32     %r21, %r15, %r12;
    add.s32     %r12, %r5, %r21;
    .loc    17  16  0
 //  16         num_inliers += j*(j+7);
    add.s32     %r22, %r15, %r12;
    add.s32     %r12, %r6, %r22;
    .loc    17  17  0
 //  17         num_inliers += j*(j+8);
    add.s32     %r23, %r15, %r12;
    add.s32     %r12, %r7, %r23;
    .loc    17  18  0
 //  18         num_inliers += j*(j+9);
    add.s32     %r24, %r15, %r12;
    add.s32     %r12, %r8, %r24;
    .loc    17  19  0
 //  19         num_inliers += j*(j+10);
    add.s32     %r25, %r15, %r12;
    add.s32     %r12, %r9, %r25;
    .loc    17  20  0
 //  20         num_inliers += j*(j+11);
    add.s32     %r26, %r15, %r12;
    add.s32     %r12, %r10, %r26;
    .loc    17  21  0
 //  21         num_inliers += j*(j+12);
    add.s32     %r27, %r15, %r12;
    add.s32     %r12, %r11, %r27;
    add.s32     %r14, %r14, 1;
    add.s32     %r11, %r11, 12;
    add.s32     %r10, %r10, 11;
    add.s32     %r9, %r9, 10;
    add.s32     %r8, %r8, 9;
    add.s32     %r7, %r7, 8;
    add.s32     %r6, %r6, 7;
    add.s32     %r5, %r5, 6;
    add.s32     %r4, %r4, 5;
    add.s32     %r3, %r3, 4;
    add.s32     %r2, %r2, 3;
    add.s32     %r1, %r1, 2;
    mov.u32     %r28, 1600;
    setp.ne.s32     %p1, %r1, %r28;
    @%p1 bra    $Lt_0_2562;
    cvt.u32.u16     %r29, %tid.x;
    mov.u32     %r30, -1;
    setp.ne.u32     %p2, %r29, %r30;
    @%p2 bra    $Lt_0_3074;
    .loc    17  25  0
 //  22     }
 //  23  
 //  24     if (threadIdx.x == -1)
 //  25         d_out[threadIdx.x] = num_inliers;
    cvt.rn.f32.s32  %f1, %r12;
    ld.param.u32    %r31, [__cudaparm__Z12less_threadsPf_d_out];
    mul24.lo.u32    %r32, %r29, 4;
    add.u32     %r33, %r31, %r32;
    st.global.f32   [%r33+0], %f1;
$Lt_0_3074:
    .loc    17  26  0
 //  26  }
    exit;
$LDWend__Z12less_threadsPf:
    } // _Z12less_threadsPf

    .entry _Z12more_threadsPf (
        .param .u32 __cudaparm__Z12more_threadsPf_d_out)
    {
    .reg .u32 %r<19>;
    .reg .f32 %f<3>;
    .reg .pred %p<4>;
    .loc    17  28  0
 //  27  
 //  28  __global__ void more_threads(float *d_out) {
$LBB1__Z12more_threadsPf:
    mov.s32     %r1, 0;
    mov.s32     %r2, 0;
    mov.s32     %r3, 0;
    mov.s32     %r4, %r5;
    mov.s32     %r6, 0;
$Lt_1_2562:
 //<loop> Loop body line 28, nesting depth: 1, iterations: 800
    .loc    17  32  0
 //  29     int num_inliers;
 //  30     for (int j=0;j<800;++j) {
 //  31         // Do 4 computations
 //  32         num_inliers += j*(j+1);
    mul.lo.s32  %r7, %r6, %r6;
    add.s32     %r8, %r4, %r6;
    add.s32     %r4, %r7, %r8;
    .loc    17  33  0
 //  33         num_inliers += j*(j+2);
    add.s32     %r9, %r7, %r4;
    add.s32     %r4, %r1, %r9;
    .loc    17  34  0
 //  34         num_inliers += j*(j+3);
    add.s32     %r10, %r7, %r4;
    add.s32     %r4, %r2, %r10;
    .loc    17  35  0
 //  35         num_inliers += j*(j+4);
    add.s32     %r11, %r7, %r4;
    add.s32     %r4, %r3, %r11;
    add.s32     %r6, %r6, 1;
    add.s32     %r3, %r3, 4;
    add.s32     %r2, %r2, 3;
    add.s32     %r1, %r1, 2;
    mov.u32     %r12, 1600;
    setp.ne.s32     %p1, %r1, %r12;
    @%p1 bra    $Lt_1_2562;
    cvt.u32.u16     %r13, %tid.x;
    mov.u32     %r14, -1;
    setp.ne.u32     %p2, %r13, %r14;
    @%p2 bra    $Lt_1_3074;
    .loc    17  38  0
 //  36     }
 //  37     if (threadIdx.x == -1)
 //  38         d_out[threadIdx.x] = num_inliers;
    cvt.rn.f32.s32  %f1, %r4;
    ld.param.u32    %r15, [__cudaparm__Z12more_threadsPf_d_out];
    mul24.lo.u32    %r16, %r13, 4;
    add.u32     %r17, %r15, %r16;
    st.global.f32   [%r17+0], %f1;
$Lt_1_3074:
    .loc    17  39  0
 //  39  }
    exit;
$LDWend__Z12more_threadsPf:
    } // _Z12more_threadsPf

Обратите внимание, что оба ядра должны выполнять в целом одинаковый объем работы (если threadIdx.x == -1 - это хитрость, которая не позволяет компилятору оптимизировать все и оставить пустое ядро). Работа должна быть такой же, как more_threads использует в 4 раза больше потоков, но каждый поток выполняет в 4 раза меньше работы.

Значимые результаты от профилировщика следующие: L: 1011 *

more_threads: время выполнения GPU = 1474 мкс, рег. На поток = 6, занятость = 1, ветвь = 83746, divergent_branch = 26, инструкции = 584065, запрос gst = 1084552

less_threads: время выполнения GPU = 921 мкс, рег. На поток = 14, занятость = 0,25, ветвь = 20956, divergent_branch = 26, инструкции = 312663, запрос gst = 677381

Как я уже говорил ранее, время работы ядра с использованием большего количества потоков больше, это может быть связано с увеличением количества инструкций.

Почему есть еще инструкции?

Почему существует какое-либо ветвление, , не говоря уже о расходящемся ветвлении, учитывая отсутствие условного кода?

Почему существуют gst запросы , когда нет доступа к глобальной памяти?

Что здесь происходит!

Спасибо

Обновление

Добавлен код PTX и исправлен CUDA C, поэтому он должен компилироваться

Ответы [ 3 ]

4 голосов
/ 16 марта 2010

Поскольку в вашем коде есть только арифметические инструкции, вам не нужно очень много людей, чтобы скрыть задержку арифметических единиц. Действительно, даже если у вас есть инструкции по памяти, вы можете максимизировать производительность, занимая ~ 50% при условии, что ваши операции чтения / записи эффективны. См. Записанную Advanced CUDA C презентацию для получения дополнительной информации о занятости и производительности.

В вашем случае, учитывая, что вашему ядру не требуется высокая занятость для насыщения арифметических единиц, вы будете иметь лучшую производительность, используя меньше блоков большего размера, чем блоков меньшего размера, так как затраты на запуск блоков высоки. Однако в целом стоимость запуска блоков незначительна по сравнению со временем фактического запуска кода.

Почему есть еще инструкции?

Помните, что счетчики рассчитывают не на блок (он же CTA), а на каждый SM (потоковый мультипроцессор) или TPC (Texture Processing Cluster), который представляет собой группу из двух или трех SM в зависимости от вашего устройства. Количество инструкций по SM.

Справедливо ожидать, что ядро ​​less_threads будет иметь меньше инструкций, однако вы запускаете в четыре раза больше перекосов в блоке, что означает, что каждый SM выполнит код примерно в четыре раза больше раз. Учитывая более короткий код ядра, ваши измерения не кажутся необоснованными.

Почему существует разветвление?

На самом деле у вас есть условный код: 

for (int j=0;j<800;++j)

Это имеет условие, однако все потоки внутри деформации действительно выполняют один и тот же путь, поэтому он не расходится. Я думаю, что расхождение в коде администрирования где-то, вы могли бы взглянуть на код PTX, чтобы проанализировать это, если вы беспокоитесь. 26 очень мало по сравнению с количеством выполненных инструкций, поэтому это не повлияет на вашу производительность.

Почему есть gst запросы?

В вашем коде у вас есть: 

if (threadIdx.x == -1)
  d_out[blockIdx.x*blockDim.x+threadIdx.x] = num_inliers;

Это будет обработано модулем загрузки / хранения и, следовательно, засчитано, даже если это не приведет к фактической транзакции. Счетчики gst_32 / gst_64 / gst_128 указывают фактическую передачу памяти (ваше устройство имеет вычислительные возможности 1.2 или 1.3, старые устройства имеют разные наборы счетчиков).

4 голосов
/ 18 марта 2010

Две функции не выполняют одинаковое количество работы.

more_threads<<<780, 128>>>():

  • 780 блоков
  • 128 потоков на блок
  • 4 мул на цикл
  • 8 добавлений за цикл
  • 780 * 128 * 800 * (4 + 8) = 958 464 000 флопов

less_threads<<<780, 32>>>():

  • 780 блоков
  • 32 потока на блок
  • 12 муль за цикл
  • 24 добавления за цикл
  • 780 * 32 * 800 * (12 + 24) = 718 848 000 флопов

Итак, more_threads выполняет больше работы, чем меньше потоков, поэтому количество инструкций увеличивается, а more_threads медленнее. Чтобы исправить more_threads, сделайте только 3 вычисления внутри цикла: 780 * 128 * 800 * (3 + 6) = 718 848 000.

0 голосов
/ 16 марта 2010

  1. две функции имеют разное количество строк кода, поэтому разное количество инструкций

  2. для цикла реализован с использованием веток. последняя строка кода всегда расходится

  3. запрос глобального хранилища не совпадает с глобальным счетом. операция настроена, но не зафиксирована.

Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.

Нет похожих вопросов

...