Должен ли тип thread_block передаваться по ссылке?

Question

Вопрос

При прохожденииthread_groupвведите объекты в функцию устройства, есть ли предпочтение передаче по ссылке, а не передаче по значению?

• Один из них «правильный»?
• В чем различия каждого подхода
• Когда следует отдать предпочтение каждому подходу

Примеры

Подобные примеры в руководстве по программированию и блоге разработчиков , похоже, решают эту проблему по-разному.

Руководство по программированию

      __device__
int sum(const thread_block& g, int *x, int n) {
    // ...
    g.sync()
    return total;
}

Блог разработчиков

      __device__
int sum(thread_block block, int *x, int n) {
    ...
    block.sync();
    ...
    return total;
}

Дополнительная информация

В руководстве по программированию также говорится о построении неявных групп:

Хотя вы можете создать неявную группу в любом месте кода, это опасно. Создание дескриптора неявной группы — это коллективная операция: в ней должны участвовать все потоки в группе. Если группа была создана в условной ветке, до которой доходят не все потоки, это может привести к взаимоблокировкам или повреждению данных. По этой причине рекомендуется заранее создать дескриптор неявной группы (как можно раньше, до того, как произойдет какое-либо ветвление) и использовать этот дескриптор во всем ядре. Дескрипторы групп должны быть инициализированы во время объявления (нет конструктора по умолчанию) по той же причине, и их копирование не рекомендуется.

Это заставило бы меня поверить, что передача их по ссылке предпочтительнее, но я признаю, что в основе различных кооперативных групп более чем достаточно подробностей, поэтому, вероятно, я упустил некоторые нюансы. Будет ли передача по значению считаться «конструированием копирования» и поэтому не будет поощряться?

Я не заметил никакой разницы в производительности или результатах при использовании любого из них, но, возможно, я просто не проверил правильный крайний случай; или «неопределенное поведение» может просто работать таким образом, чтобы не вызывать проблем.

Answer 1

Сначала несколько наблюдений:

1. Сообщение в блоге, на которое вы ссылаетесь, относится к 2017 году, когда была представлена ​​предварительная версия этой функции, документация актуальна. Уже на этом основании вам следует отдать предпочтение идиоме const с передачей по ссылке, поскольку исходный код новее.
2. Как вы сами доказали, поскольку CUDA использует сильно урезанную реализацию объектной модели C++, а компилятор любит расширение встроенных функций для повышения производительности, очень маловероятно, что вы встретите в реальном мире случаи, когда компилятор будет генерировать другой код для два случая.

В результате я считаю, что вам следует использовать версию const с передачей по ссылке, как с точки зрения корректности языка C++, так и потому, что текущая документация предполагает, что вам следует это делать. Вероятно, есть крайние случаи, когда кто-то где-то и когда-то обжигался из-за построения копирования в версии с передачей по значению, но я подозреваю, что вам придется очень постараться, чтобы это произошло. Caveat emptor и все такое….