Предотвращение расхождения потоков в ядре CUDA с условными присваиваниями

Question

Я пытаюсь найти способы избежать расхождения потоков ( ветвление или деформация деформации ) в моем ядре CUDA.

ДляНапример, у меня есть следующее условное присвоение (a и b являются char значениями, x и y являются unsigned int значениями):

if (a == b) { ++x; }
else        { ++y; }

Или, альтернативно:

if (a == b) { ++x; }
if (a != b) { ++y; }

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

Я смотрел в типах приведения типов , но приведения нетbool до int.Я думаю, что может быть какой-то трюк с min, max и абсолютными значениями (например, __ sad ), чтобы получить соответствующий целочисленный результат, который нужно добавить для каждого случая (то есть 1, 0 или 0, 1).

Кажется, что не существует обычной функции абсолютного значения int, но я вижу следующее:

Рассчитать |х - у |+ z, сумма абсолютной разности.

__device__ ​ unsigned int __sad ( int  x, int  y, unsigned int  z )

Который, я полагаю, мог бы предоставить аргумент z = 0, чтобы получить нормальное абсолютное значение.Может быть, что-то вроде:

const unsigned int mu = __sad(a, b, 1);
const unsigned int mv = __sad(a, b, 0);
const int u = __nv_min(1, mu);
const int v = __nv_min(1, mv);
x += u;
y += v;

Однако нет функции min (см. связанный вопрос ).

Answer 1

tl; dr: попробуйте избежать таких предполагаемых микрооптимизаций.

Давайте посмотрим, сможем ли мы определить, какие есть отличия (если таковые имеются) от первоначальной формулировки, предложенной в вопросе:

if (a == b) { ++x; }
else        { ++y; }

и формулировка, предложенная в другой ответ :

x += (a == b);
y += (a != b);

мы будем использовать этот тестовый код:

$ cat t1513.cu
__global__ void k(char a, char b, unsigned int *dx, unsigned int *dy){

  unsigned int x=*dx;
  unsigned int y=*dy;
#ifndef USE_OPT
  if (a == b)
{
    ++x;
} else {
    ++y;
}
#else
x += (a == b);
y += (a != b);
#endif

  *dy = y;
  *dx = x;
}


$ nvcc -c t1513.cu
$ cuobjdump -sass t1513.o >out1.sass
$ nvcc -c t1513.cu -DUSE_OPT
$ cuobjdump -sass t1513.o >out2.sass
$ diff out1.sass out2.sass
28,29c28,29
<         /*0078*/                   BFE R7, R7, 0x1000;          /* 0x7000c0400071dc23 */
<                                                                 /* 0x22e04283f2828287 */
---
>         /*0078*/                   BFE R9, R7, 0x1000;          /* 0x7000c04000725c23 */
>                                                                 /* 0x22804283f2804287 */
31,41c31,41
<         /*0090*/                   ISET.EQ.AND R7, R8, R7, PT;  /* 0x110e00001c81dc23 */
<         /*0098*/                   LOP32I.AND R7, R7, 0x1;      /* 0x380000000471dc02 */
<         /*00a0*/                   LOP32I.XOR R8, R7, 0x1;      /* 0x3800000004721c82 */
<         /*00a8*/                   IADD R8, R6, R8;             /* 0x4800000020621c03 */
<         /*00b0*/                   IADD R7, R0, R7;             /* 0x480000001c01dc03 */
<         /*00b8*/                   ST.E [R4], R8;               /* 0x9400000000421c85 */
<                                                                 /* 0x200000000002f047 */
<         /*00c8*/                   ST.E [R2], R7;               /* 0x940000000021dc85 */
<         /*00d0*/                   EXIT;                        /* 0x8000000000001de7 */
<         /*00d8*/                   BRA 0xd8;                    /* 0x4003ffffe0001de7 */
<         /*00e0*/                   NOP;                         /* 0x4000000000001de4 */
---
>         /*0090*/                   ISET.NE.AND R7, R8, R9, PT;  /* 0x128e00002481dc23 */
>         /*0098*/                   ISET.EQ.AND R8, R8, R9, PT;  /* 0x110e000024821c23 */
>         /*00a0*/                   LOP32I.AND R7, R7, 0x1;      /* 0x380000000471dc02 */
>         /*00a8*/                   IADD R7, R6, R7;             /* 0x480000001c61dc03 */
>         /*00b0*/                   LOP32I.AND R6, R8, 0x1;      /* 0x3800000004819c02 */
>         /*00b8*/                   ST.E [R4], R7;               /* 0x940000000041dc85 */
>                                                                 /* 0x2000000002f04287 */
>         /*00c8*/                   IADD R6, R0, R6;             /* 0x4800000018019c03 */
>         /*00d0*/                   ST.E [R2], R6;               /* 0x9400000000219c85 */
>         /*00d8*/                   EXIT;                        /* 0x8000000000001de7 */
>         /*00e0*/                   BRA 0xe0;                    /* 0x4003ffffe0001de7 */
$

Изучая приведенный выше вывод diff, мысм .:

1. В любой реализации нет разветвления (и даже не предиката).
2. Предположительно «оптимизированный» случай почти идентичен, за исключением того, что он на 1 инструкцию длиннеечем в случае if / else.

Да, я понимаю, что это не "ваш код".Я могу работать только с тем, что представлено.

Это дает мне представление о том, что следующие типы преобразований:

1. Требуют усилий (потенциально потерянное время)
2. Не можетможет привести к улучшению производительности
3. Может запутать код, усложняя обслуживание

Продолжайте, конечно, как хотите.

Answer 2

Как отметил полезный комментарий, я переосмыслил проблему.Следующее работает и использует простое преобразование bool в int :

x += (a == b);
y += (a != b);

Изучение файла сборки PTX до и после этого изменения (несколько мест в ядре), количество ветвейбыл сокращен с 39 до 9, так что это внесло значительные изменения.Компилятор nvcc не оптимизировал их сам по себе, особенно в тех случаях, когда были операторы if / then / else глубиной два или три уровня, например:

bool ag = (ca == '.');
bool bg = (cb == '.');

bool agx = ag && apg;
bool bgx = bg && bpg;
bool gx = agx || bgx;

if (ag || bg)
{
    if (ag && bg)
    {
        // ignore
    } else {
        if (!gx)
        {
            ++gs;
            ++ps;
        }
        apg = ag;
        bpg = bg;
    }
} else {
    if (ca == cb)
    {
        ++ms;
        ++ps;
    } else {
        ++ns;
        ++ps;
    }
    apg = false;
    bpg = false;
}

Как только я смог уменьшитьвсе назначения для логических выражений (здесь два из шести назначений после преобразования из исходного ядра):

apg = (apg && !!(ag && bg)) || ((ag || bg) && !(ag && bg) && ag)
bpg = (bpg && !!(ag && bg)) || ((ag || bg) && !(ag && bg) && bg)

Мне удалось упростить эти выражения:

apg  = (ag && !bg) || (ag && apg)
bpg  = (!ag && bg) || (bg && bpg)

И в двух случаях мне удалось объединить несколько выражений (несколько назначений) в одно логическое выражение.В конечном итоге полный набор условий был сокращен до:

ps += ((ca != '.') && (cb != '.')) || ((ca != '.') && !bpg) || ((cb != '.') && !apg);
ms += (ca == cb) && (ca != '.') && (cb != '.');
apg = ((ca == '.') && (cb != '.')) || ((ca == '.') && apg);
bpg = ((ca != '.') && (cb == '.')) || ((cb == '.') && bpg);

Основываясь на методе из этого ответа , я обнаружил, что число вещественных ветвей в моем ядрев конечном итоге были сокращены с 39 до 12:

cuobjdump -sass kernel_original.o > kernel_original.sass
grep BRA kernel.sass | wc -l
39

cuobjdump -sass kernel_simplified.o > kernel_simplified.sass
grep BRA kernel_opt.sass | wc -l
12

Answer 3

tl; dr: сначала рассмотрите большую картину, прежде чем применять такие предполагаемые микрооптимизации.

Глядя на пример кода Роберта, моя первая мысль была

++*( (a==b) ? &x : &y);

Однако я был намой мобильный телефон и сам не смог проверить разборку этого.

Роберт был достаточно любезен, чтобы вставить его в свое тестовое ядро, и опубликовал разность SASS этой идеи по сравнению с исходным кодом if / else, размещенным в вопросе.:

$ cuobjdump -sass t1513.o >out3.sass
$ diff out1.sass out3.sass
13,44c13,52
<                                                                 /* 0x2230427042004307 */
<         /*0008*/                   MOV R1, c[0x0][0x44];        /* 0x2800400110005de4 */
<         /*0010*/                   MOV R4, c[0x0][0x150];       /* 0x2800400540011de4 */
<         /*0018*/                   MOV R5, c[0x0][0x154];       /* 0x2800400550015de4 */
<         /*0020*/                   MOV R2, c[0x0][0x148];       /* 0x2800400520009de4 */
<         /*0028*/                   MOV R3, c[0x0][0x14c];       /* 0x280040053000dde4 */
<         /*0030*/                   LD.E R6, [R4];               /* 0x8400000000419c85 */
<         /*0038*/                   LDC.U8 R7, c[0x0][0x141];    /* 0x1400000507f1dc06 */
<                                                                 /* 0x2272028042824047 */
<         /*0048*/                   LD.E R0, [R2];               /* 0x8400000000201c85 */
<         /*0050*/                   LDC.U8 R8, c[0x0][0x140];    /* 0x1400000503f21c06 */
<         /*0058*/                   I2I.S16.S8 R7, R7;           /* 0x1c0000001c11de84 */
<         /*0060*/                   I2I.S16.S8 R8, R8;           /* 0x1c00000020121e84 */
<         /*0068*/                   LOP32I.AND R7, R7, 0xff;     /* 0x38000003fc71dc02 */
<         /*0070*/                   LOP32I.AND R8, R8, 0xff;     /* 0x38000003fc821c02 */
<         /*0078*/                   BFE R7, R7, 0x1000;          /* 0x7000c0400071dc23 */
<                                                                 /* 0x22e04283f2828287 */
<         /*0088*/                   BFE R8, R8, 0x1000;          /* 0x7000c04000821c23 */
<         /*0090*/                   ISET.EQ.AND R7, R8, R7, PT;  /* 0x110e00001c81dc23 */
<         /*0098*/                   LOP32I.AND R7, R7, 0x1;      /* 0x380000000471dc02 */
<         /*00a0*/                   LOP32I.XOR R8, R7, 0x1;      /* 0x3800000004721c82 */
<         /*00a8*/                   IADD R8, R6, R8;             /* 0x4800000020621c03 */
<         /*00b0*/                   IADD R7, R0, R7;             /* 0x480000001c01dc03 */
<         /*00b8*/                   ST.E [R4], R8;               /* 0x9400000000421c85 */
<                                                                 /* 0x200000000002f047 */
<         /*00c8*/                   ST.E [R2], R7;               /* 0x940000000021dc85 */
<         /*00d0*/                   EXIT;                        /* 0x8000000000001de7 */
<         /*00d8*/                   BRA 0xd8;                    /* 0x4003ffffe0001de7 */
<         /*00e0*/                   NOP;                         /* 0x4000000000001de4 */
<         /*00e8*/                   NOP;                         /* 0x4000000000001de4 */
<         /*00f0*/                   NOP;                         /* 0x4000000000001de4 */
<         /*00f8*/                   NOP;                         /* 0x4000000000001de4 */
---
>                                                                      /* 0x2270420042304307 */
>         /*0008*/                   MOV R1, c[0x0][0x44];             /* 0x2800400110005de4 */
>         /*0010*/                   MOV R10, c[0x0][0x148];           /* 0x2800400520029de4 */
>         /*0018*/                   IADD32I R1, R1, -0x8;             /* 0x0bffffffe0105c02 */
>         /*0020*/                   MOV R11, c[0x0][0x14c];           /* 0x280040053002dde4 */
>         /*0028*/                   LDC.U8 R0, c[0x0][0x141];         /* 0x1400000507f01c06 */
>         /*0030*/                   MOV R8, c[0x0][0x150];            /* 0x2800400540021de4 */
>         /*0038*/                   MOV R9, c[0x0][0x154];            /* 0x2800400550025de4 */
>                                                                      /* 0x2232423240423047 */
>         /*0048*/                   LD.E R4, [R10];                   /* 0x8400000000a11c85 */
>         /*0050*/                   I2I.S16.S8 R0, R0;                /* 0x1c00000000101e84 */
>         /*0058*/                   LD.E R5, [R8];                    /* 0x8400000000815c85 */
>         /*0060*/                   LDC.U8 R2, c[0x0][0x140];         /* 0x1400000503f09c06 */
>         /*0068*/                   LOP32I.AND R0, R0, 0xff;          /* 0x38000003fc001c02 */
>         /*0070*/                   I2I.S16.S8 R2, R2;                /* 0x1c00000008109e84 */
>         /*0078*/                   BFE R0, R0, 0x1000;               /* 0x7000c04000001c23 */
>                                                                      /* 0x2283f282b2028287 */
>         /*0088*/                   LOP32I.AND R2, R2, 0xff;          /* 0x38000003fc209c02 */
>         /*0090*/                   BFE R3, R2, 0x1000;               /* 0x7000c0400020dc23 */
>         /*0098*/                   ISETP.NE.AND P0, PT, R3, R0, PT;  /* 0x1a8e00000031dc23 */
>         /*00a0*/                   LOP.OR R3, R1, c[0x0][0x24];      /* 0x680040009010dc43 */
>         /*00a8*/               @P0 IADD32I R3, R3, 0x4;              /* 0x080000001030c002 */
>         /*00b0*/                   LOP32I.AND R3, R3, 0xffffff;      /* 0x3803fffffc30dc02 */
>         /*00b8*/                   SEL R0, R4, R5, !P0;              /* 0x2010000014401c04 */
>                                                                      /* 0x22f042e3f2e28047 */
>         /*00c8*/                   STL.64 [R1], R4;                  /* 0xc800000000111ca5 */
>         /*00d0*/                   IADD32I R0, R0, 0x1;              /* 0x0800000004001c02 */
>         /*00d8*/                   STL [R3], R0;                     /* 0xc800000000301c85 */
>         /*00e0*/                   LDL.64 R6, [R1];                  /* 0xc000000000119ca5 */
>         /*00e8*/                   ST.E [R8], R7;                    /* 0x940000000081dc85 */
>         /*00f0*/                   ST.E [R10], R6;                   /* 0x9400000000a19c85 */
>         /*00f8*/                   EXIT;                             /* 0x8000000000001de7 */
>         /*0100*/                   BRA 0x100;                        /* 0x4003ffffe0001de7 */
>         /*0108*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0110*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0118*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0120*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0128*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0130*/                   NOP;                              /* 0x4000000000001de4 */
>         /*0138*/                   NOP;                              /* 0x4000000000001de4 */
$

Роберт пришел к выводу, что компилятор решил использовать предикацию в этом случае.

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

Однако я подумал о том, чтобы просто заменить все тело контрольного примера на мою ++*( (a==b) ? dx : dy); однострочную, чтопривели к лучшему виду сборки:

        /*0008*/                   MOV R1, c[0x0][0x44];             /* 0x2800400110005de4 */
        /*0010*/                   LDC.U8 R0, c[0x0][0x141];         /* 0x1400000507f01c06 */
        /*0018*/                   LDC.U8 R2, c[0x0][0x140];         /* 0x1400000503f09c06 */
        /*0020*/                   I2I.S16.S8 R0, R0;                /* 0x1c00000000101e84 */
        /*0028*/                   I2I.S16.S8 R2, R2;                /* 0x1c00000008109e84 */
        /*0030*/                   LOP32I.AND R0, R0, 0xff;          /* 0x38000003fc001c02 */
        /*0038*/                   LOP32I.AND R2, R2, 0xff;          /* 0x38000003fc209c02 */
                                                                     /* 0x228202c042804237 */
        /*0048*/                   BFE R0, R0, 0x1000;               /* 0x7000c04000001c23 */
        /*0050*/                   BFE R3, R2, 0x1000;               /* 0x7000c0400020dc23 */
        /*0058*/                   MOV R2, c[0x0][0x148];            /* 0x2800400520009de4 */
        /*0060*/                   ISETP.NE.AND P0, PT, R3, R0, PT;  /* 0x1a8e00000031dc23 */
        /*0068*/                   MOV R0, c[0x0][0x14c];            /* 0x2800400530001de4 */
        /*0070*/                   SEL R2, R2, c[0x0][0x150], !P0;   /* 0x2010400540209c04 */
        /*0078*/                   SEL R3, R0, c[0x0][0x154], !P0;   /* 0x201040055000dc04 */
                                                                     /* 0x20000002f04283f7 */
        /*0088*/                   LD.E R0, [R2];                    /* 0x8400000000201c85 */
        /*0090*/                   IADD32I R4, R0, 0x1;              /* 0x0800000004011c02 */
        /*0098*/                   ST.E [R2], R4;                    /* 0x9400000000211c85 */
        /*00a0*/                   EXIT;                             /* 0x8000000000001de7 */
        /*00a8*/                   BRA 0xa8;                         /* 0x4003ffffe0001de7 */
        /*00b0*/                   NOP;                              /* 0x4000000000001de4 */
        /*00b8*/                   NOP;                              /* 0x4000000000001de4 */

Этот код выглядит для меня лучше, чем тестовый сценарий Роберта (сам по себе).Но это, вероятно, бесполезно для vallismortis, потому что в его случае переменные не будут в адресуемой памяти.

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