Кооперативные группы
Современный API для гибких областей синхронизации.
«Кооперативные группы» — бесплатный урок CUDA Academy на CoddyKit. Это урок 4 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.
Более удобный интерфейс синхронизации
Необработанные маски и встроенные функции работают, но с ними легко ошибиться. Кооперативные группы объединяют потоки в объекты, которым можно задавать имя и размер и которые можно явно синхронизировать.
#include <cooperative_groups.h>
namespace cg = cooperative_groups;Получите весь блок
Начните с получения дескриптора своего блока потоков. this_thread_block возвращает группу, которую можно синхронизировать так же, как с помощью syncthreads, но в виде объекта.
cg::thread_block block = cg::this_thread_block();Синхронизация через группу
Вызов sync для блока — это современный барьер. Он делает ровно то же, что и syncthreads, но понятнее показывает, что это метод нужной вам группы.
block.sync();Выделите фрагмент размера варпа
Блок можно разделить на фрагменты фиксированного размера — tiles. Фрагмент из 32 линий даёт группу размером с варп, где есть понятные методы вместо необработанных масок перемешивания.
auto warp = cg::tiled_partition<32>(block);Методы заменяют маски
Фрагментированная группа предоставляет shfl_down и другие операции без аргумента-маски. Группа уже знает своих участников, поэтому интерфейс остаётся коротким и безопасным.
val += warp.shfl_down(val, offset);Знайте свою позицию
Каждая группа предоставляет сведения о своих участниках и вашей позиции. thread_rank возвращает ваш индекс внутри группы, а size возвращает количество содержащихся в ней потоков.
int rank = warp.thread_rank();Фрагменты могут быть и меньше
Фрагменты не обязаны иметь ширину 32. tiled_partition размера 8 или 16 создаёт подварповые группы, удобные, когда ваши данные естественным образом объединяются в небольшие наборы.
Редукции на уровне группы
Библиотека предоставляет готовые коллективные операции. Группа reduce суммирует фрагмент одним вызовом, скрывая цикл по смещениям, который раньше вы писали вручную.
int total = cg::reduce(warp, val, cg::plus<int>());Сетки с синхронизацией
Самая масштабная группа — это группа сетки. При кооперативном запуске каждый блок может синхронизироваться на одном барьере, чего обычное ядро делать не может.
Требуется кооперативный запуск
Синхронизация всей сетки работает только при запуске ядра с помощью cudaLaunchCooperativeKernel и при поддержке этой возможности GPU. Обычный запуск её не разрешает.
Зачем это нужно
Кооперативные группы делают код варпа понятным и переносимым: не нужны маски, которыми приходится управлять вручную, области действия ясны, а повторно используемые коллективные операции соответствуют вашему замыслу.
Быстрая проверка
Вспомните, как создать кооперативную группу размера варпа из блока.
Итоги
Кооперативные группы превращают маски в именованные объекты: разделяют блок на фрагменты, вызывают reduce и даже синхронизируют целую сетку. Теперь вы управляете CUDA на уровне варпа. ✨
Изучай C++ с ИИ-репетитором — бесплатно
Пиши и запускай код прямо в браузере, получай мгновенную помощь от ИИ-репетитора 24/7 и продолжи учиться на сайте или в приложении.
- Курсы
- 30
- Уроки
- 120
Часто задаваемые вопросы
Урок «Кооперативные группы» бесплатный?
Да — полный текст урока «Кооперативные группы» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.
Чему я научусь в уроке «Кооперативные группы»?
Современный API для гибких областей синхронизации. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.
Нужен ли мне опыт, чтобы начать CUDA Academy?
Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 4 из 4.
Сколько времени занимает урок «Кооперативные группы»?
Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.
Можно ли писать и запускать код в этом уроке CUDA Academy?
Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.
Все уроки этого курса
- Варпы, lanes и маски
- __shfl_down_sync для свёрток
- Функции голосования и ballot
- Кооперативные группы