#c #c #optimization #cuda #gpgpu
#c #c #оптимизация #cuda #gpgpu
Вопрос:
У меня есть 2 очень похожие функции ядра, в том смысле, что код почти одинаков, но с небольшим отличием. В настоящее время у меня есть 2 варианта:
- Напишите 2 разных метода (но очень похожих)
- Напишите одно ядро и поместите блоки кода, которые отличаются, в инструкцию if / else
Насколько сильно оператор if повлияет на производительность моего алгоритма?
Я знаю, что ветвления нет, поскольку все потоки во всех блоках будут вводить либо if, либо else.
Итак, снизит ли одна инструкция if мою производительность, если функция ядра вызывается много раз?
Комментарии:
1. Почему бы вам не попробовать оба и не рассчитать время?
Ответ №1:
У вас есть третья альтернатива, которая заключается в использовании шаблонов C и превращении переменной, которая используется в инструкции if / switch, в параметр шаблона. Создайте экземпляр каждой версии ядра, которое вам нужно, и тогда у вас будет несколько ядер, выполняющих разные задачи, без расхождений ветвей или условной оценки, о которых нужно беспокоиться, потому что компилятор оптимизирует мертвый код и ветвление вместе с ним.
Возможно, что-то вроде этого:
template<int action>
__global__ void kernel()
{
switch(action) {
case 1:
// First code
break;
case 2:
// Second code
break;
}
}
template void kernel<1>();
template void kernel<2>();
Комментарии:
1. @talonmies… Этот ответ — самый полезный, который я когда-либо видел на SO. Количество ядер CUDA в моем коде сократилось примерно в 4 раза. 🙂
2. шаблоны очень полезны для передачи размера блока потока. Это делает размер блока статичным и в то же время адаптивным по отношению к конкретным вычислительным возможностям GPU. (Нет способа использовать директиву #define из C таким же образом.) Смотрите matrixmul.cu пример.
3. Это по-прежнему не работает с переменными в качестве аргумента шаблона, но вы можете обойти это с помощью сложных инструкций if…
kernel<a><<<1,1>>>(1,2);не работает, ноif (a == 1) { kernel<1><<<1,1>>>(1,2)} работает.4. @XapaJIaMnu: Насколько мне известно, аргументами шаблона могут быть только параметры типа или постоянные выражения. Что вы подразумеваете под «переменными в качестве аргументов шаблона»?
5. @talonmies, упс, я этого не знал. В принципе, я хотел использовать переменную в шаблоне и ограничить запуск этого шаблона только в том случае, если переменная является частью предварительно созданных шаблонов, и выдавать ошибку времени выполнения в противном случае. Однако после прочтения кажется, что это невозможно с C . В примере из моего предыдущего комментария я хочу, чтобы код создавал исключение во время выполнения, если
a != 1иa != 2.
Ответ №2:
Это немного снизит вашу производительность, особенно если это происходит во внутреннем цикле, поскольку вы время от времени тратите впустую слот для выдачи команд, но это далеко не так сильно, как если бы деформация была расходящейся.
Однако, если это имеет большое значение, возможно, стоит перенести условие за пределы цикла. Однако, если деформация действительно расходится, подумайте о том, как удалить ветвление: например, вместо
if (i>0) {
x = 3;
} else {
x = y;
}
попробуйте
x = ((i>0)*3) | ((i<3)*y);