Самый чистый способ избежать ветвления с помощью шаблонов

У меня есть ядро ​​​​cuda, которое принимает некоторые аргументы шаблона bool, которые помогают явно оптимизировать неиспользуемые разделы ядра. Мне интересно, есть ли более чистый/компактный способ создания экземпляра вызова ядра, чем мое решение ниже.

Шаблон устройства

template <bool pred1, bool pred2>
__global__ void foo(args)
{
  /* probably doesn't need to exist, but just to improve readability in the kernel */
  constexpr bool predAll = pred1 && pred2;

  if (predAll || pred1 || pred2 || someVariablePred) {
  /* working hard */
  } else {
  /* or hardly working */
  }
}

Код рабочего хоста

const bool pred1 = variableCond1 == variableCond2, pred2 = variableCond3 == variableCond4;

/* working, compiles, but cumbersome */
for (int i = 0; i < bigNumber; ++i) {
  if (pred1) {
    if (pred2) {
      foo<true, true><<<dimGrid, dimBlock>>>(args...);
    }
    foo<true, false><<<dimGrid, dimBlock>>>(args...);
  } else if (pred2) {
    foo<false, true><<<dimGrid, dimBlock>>>(args...);
  } else {
    foo<false, false><<<dimGrid, dimBlock>>>(args...);
  }
}

Хотите что-то чистое, подобное этому

/* doesn't compile */
const bool pred1 = variableCond1 == variableCond2, pred2 = variableCond3 == variableCond4;

for (int i = 0; i < bigNumber; ++i) {
  foo<pred1, pred2><<<dimGrid, dimBlock>>>(args...);
}

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


person Jacob Faib    schedule 06.11.2020    source источник
comment
Нет, шаблоны — это концепция времени компиляции (они не существуют во время выполнения), и компилятор не делает за вас никакой магии, например, генерирует все возможные комбинации и выбирает правильную во время выполнения.   -  person churill    schedule 06.11.2020
comment
вам нужно каким-то образом сопоставить условия времени выполнения с параметрами шаблона времени компиляции. Ваш рабочий код уже достаточно ясен и читабелен.   -  person 463035818_is_not_a_number    schedule 06.11.2020
comment
Еще одно не по читабельности: шаблон фактически уменьшил его, имхо. Если это близко к вашему реальному коду, то foo делает одно из двух. Вы можете закрыть это с помощью one if/else. Но таким образом у вас есть if/else внутри foo и 4 условий для выбора правильного шаблона. Так... оно того стоит? Кроме того, да, predAll = pred1 && pred2 совершенно избыточен и сбивает с толку.   -  person churill    schedule 06.11.2020
comment
К сожалению, настоящий foo значительно сложнее. pred и друзья предоставляют ярлыки через пару уровней ветвления внутри ядра, большинство из которых вызывает различные функции устройства, все из которых также принимают pred и друзей в качестве параметров шаблона.   -  person Jacob Faib    schedule 06.11.2020
comment
Я бы сказал, что цепочка if-else в коде рабочего хоста должна быть инкапсулирована в отдельную функцию, но в остальном у вас уже есть чистый подход.   -  person JaMiT    schedule 06.11.2020
comment
помню похожие вопросы, но больше не нахожу. Что вы можете сделать, так это уменьшить сложность отображения с 2x2 (для true/false x true/false) до 2 + 2.   -  person 463035818_is_not_a_number    schedule 06.11.2020
comment
Я бы поместил ваш исходный код в Godbolt, так как у меня есть ощущение, что компилятор может оптимизировать это больше, чем вы ожидаете. Например, эти условия будут перемещены за пределы замкнутых контуров. Имхо, ты лаешь не под тем деревом. Вы сделали профилирование, прежде чем обвинять этот код?   -  person Marek R    schedule 06.11.2020
comment
Да, профилирование — это то, что привело меня к созданию шаблона для начала. Чтобы обеспечить немного больше контекста, основное ядро ​​​​(foo) запускает очень жесткий цикл с привязкой к памяти. Аргументы шаблона гарантируют ядрам (если они истинны), что данные кратны некоторому числу и, следовательно, правильно выровнены в памяти для векторизованных загрузок/сохранений. Без аргументов шаблона ядрам пришлось бы вычислять на лету дорогостоящие переменные по модулю.   -  person Jacob Faib    schedule 06.11.2020


Ответы (1)


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

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

#include <iostream>
#include <string>

template <bool pred0, bool pred1, bool pred2, bool pred3, bool pred4>
void Func1()
{
  std::cout << "func1: " << pred0 << pred1 << pred2 << pred3 << pred4 << std::endl;
}

template <bool pred0, bool pred1, bool pred2, bool pred3, bool pred4>
struct Func1Wrapper
{
  static constexpr auto Ptr = Func1<pred0, pred1, pred2, pred3, pred4>;
  static void Call()
  {
    Func1<pred0, pred1, pred2, pred3, pred4>();
  }
};

template <bool pred0, bool pred1, bool pred2>
void Func2()
{
  std::cout << "func2: " << pred0 << pred1 << pred2 << std::endl;
}

template <bool pred0, bool pred1, bool pred2>
struct Func2Wrapper
{
  static constexpr auto Ptr = Func2<pred0, pred1, pred2>;
  static void Call()
  {
    Func2<pred0, pred1, pred2>();
  }
};

template <template<bool...> class F, bool... preds>
struct Ct
{
  template <bool newPred>
  using Expanded = Ct<F, preds..., newPred>;

  static void CallFunc()
  {
    F<preds...>::Call();
  }

  static auto GetFunc()
  {
    return F<preds...>::Ptr;
  }
};

template <class DeducedPreds, typename BoolLast>
void Rt2Ct(BoolLast predLast)
{
  if (predLast)
    DeducedPreds::template Expanded<true>::CallFunc();
  else
    DeducedPreds::template Expanded<false>::CallFunc();
}

template <class DeducedPreds, typename BoolFirst, typename ...BoolRest>
void Rt2Ct(BoolFirst predFirst, BoolRest... predRest)
{
  if (predFirst)
    Rt2Ct<typename DeducedPreds::template Expanded<true>, BoolRest...>(predRest...);
  else
    Rt2Ct<typename DeducedPreds::template Expanded<false>, BoolRest...>(predRest...);
}

template <template<bool...> class F, typename ...Bools>
void FuncCaller(Bools... preds)
{
  Rt2Ct<Ct<F>, Bools...>(preds...);
}

template <class DeducedPreds, typename BoolLast>
auto Rt2CtGet(BoolLast predLast)
{
  if (predLast)
    return DeducedPreds::template Expanded<true>::GetFunc();
  else
    return DeducedPreds::template Expanded<false>::GetFunc();
}

template <class DeducedPreds, typename BoolFirst, typename ...BoolRest>
auto Rt2CtGet(BoolFirst predFirst, BoolRest... predRest)
{
  if (predFirst)
    return Rt2CtGet<typename DeducedPreds::template Expanded<true>, BoolRest...>(predRest...);
  else
    return Rt2CtGet<typename DeducedPreds::template Expanded<false>, BoolRest...>(predRest...);
}

template <template<bool...> class F, typename ...Bools>
auto FuncGetter(Bools... preds)
{
  return Rt2CtGet<Ct<F>, Bools...>(preds...);
}

int main(int argc, char* argv[])
{
  FuncCaller<Func1Wrapper>(std::stoi(argv[1]), std::stoi(argv[2]), std::stoi(argv[3]), std::stoi(argv[4]), std::stoi(argv[5]));
  auto f2 = FuncGetter<Func2Wrapper>(std::stoi(argv[1]), std::stoi(argv[3]), std::stoi(argv[5]));
  f2();
  return 0;
}

Код требует C++14 (C++11 должен быть возможен с небольшим переписыванием) и свободен от CUDA, потому что я действительно не вижу, как вызов ядра вместо обычной функции меняет проблему. Конечно, вам придется добавить несколько параметров для передачи dimGrid и dimBlock, а также всех необходимых args.
Как уже отмечалось, вы можете поместить if-else за пределы цикла for. Вот почему я также реализовал подпрограмму FuncGetter.
Также обратите внимание, что вы можете убедиться, что аргументы шаблона Bools, BoolFirst и т. д. действительно относятся к типу bool.

Компиляция, вызов функции и вывод:

$ g++ -std=c++14 main.cpp && ./a.out 1 0 1 0 0
func1: 10100
func2: 110
person BlameTheBits    schedule 07.11.2020