CUDA Academy · Урок

Кооперативные группы

Современный API для гибких областей синхронизации.

Урок 4 из 413 шагов

«Кооперативные группы» — бесплатный урок 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 — локальная установка не требуется.

Все уроки этого курса

  1. Варпы, lanes и маски
  2. __shfl_down_sync для свёрток
  3. Функции голосования и ballot
  4. Кооперативные группы
← Назад к CUDA Academy