CUDA Academy · Leçon

Groupes coopératifs

Une API moderne pour des portées de synchronisation flexibles.

Leçon 4 sur 413 étapes

Groupes coopératifs est une leçon CUDA Academy gratuite sur CoddyKit. Ceci est la leçon 4 sur 4. Tu peux lire la leçon complète ci-dessous gratuitement — puis la pratiquer en direct dans le navigateur avec un éditeur de code intégré et un tuteur IA 24/7. Elle fait partie du parcours d'apprentissage CUDA Academy, et ta progression se synchronise sur le web et l'application CoddyKit. Le cours CUDA Academy comprend 4 leçons au total.

Une API de synchronisation plus claire

Les masques bruts et les intrinsèques fonctionnent, mais leur utilisation est délicate. Les Cooperative groups regroupent les threads dans des objets que vous pouvez nommer, dimensionner et synchroniser explicitement.

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

Récupérez tout le bloc

Commencez par obtenir une référence vers votre bloc de threads. this_thread_block renvoie un groupe que vous pouvez synchroniser comme avec syncthreads, mais sous forme d’objet.

cg::thread_block block = cg::this_thread_block();

Synchronisez via le groupe

L’appel de sync sur le bloc constitue la barrière moderne. Il fait exactement ce que fait syncthreads, mais se lit clairement comme une méthode du groupe concerné.

block.sync();

Découpez un bloc de taille warp

Vous pouvez diviser un bloc en tiles de taille fixe. Un tile de 32 lanes vous donne un groupe de la taille d’un warp, avec des méthodes claires au lieu de masques de brassage bruts.

auto warp = cg::tiled_partition<32>(block);

Les méthodes remplacent les masques

Un groupe découpé en tiles propose shfl_down et les fonctions associées sans argument de masque. Le groupe connaît déjà ses membres, ce qui permet de conserver une API concise et sûre.

val += warp.shfl_down(val, offset);

Connaissez votre position

Chaque groupe expose ses membres ainsi que votre position. thread_rank fournit votre indice dans le groupe, et size renvoie le nombre de threads qu’il contient.

int rank = warp.thread_rank();

Des tiles plus petits aussi

Les tiles ne doivent pas forcément avoir une largeur de 32. Un tiled_partition de 8 ou 16 crée des groupes plus petits qu’un warp, pratiques lorsque vos données sont naturellement regroupées en petits ensembles.

Réductions au niveau du groupe

La bibliothèque fournit des opérations collectives prêtes à l’emploi. Le reduce d’un groupe additionne un tile en un seul appel, en masquant la boucle de décalage que vous aviez écrite manuellement.

int total = cg::reduce(warp, val, cg::plus<int>());

Des grilles qui se synchronisent

Le groupe le plus ambitieux est le grid group. Avec un lancement coopératif, chaque bloc peut se synchroniser à une même barrière, ce qu’un noyau normal ne peut pas faire.

Lancement coopératif requis

La synchronisation de toute la grille ne fonctionne que si vous démarrez le noyau avec cudaLaunchCooperativeKernel et si le GPU la prend en charge. Un lancement normal ne l’autorise pas.

Pourquoi s’en soucier ?

Les groupes coopératifs rendent le code de warp readable et portable : plus besoin de gérer les masques manuellement, les portées sont claires et les opérations collectives réutilisables correspondent à votre intention.

Vérification rapide

Rappelez-vous comment créer, à partir d’un bloc, un groupe coopératif de la taille d’un warp.

Récapitulatif

Les groupes coopératifs transforment les masques en objets nommés : découpez un bloc en tiles, appelez reduce et synchronisez même une grille entière. Vous maîtrisez maintenant CUDA au niveau du warp. ✨

Gratuit pour commencer

Apprends C++ avec un tuteur IA — gratuit

Écris et exécute du vrai code dans ton navigateur, obtiens de l'aide instantanée d'un tuteur IA disponible 24h/24, et reprends là où tu t'es arrêté sur le web ou dans l'app.

Cours
30
Leçons
120

Questions Fréquemment Posées

La leçon « Groupes coopératifs » est-elle gratuite ?

Oui — le texte complet de « Groupes coopératifs » est gratuit à lire ici sur le web. Pour la pratiquer de manière interactive (un éditeur de code intégré et un tuteur IA 24/7) et déverrouiller le reste du cours CUDA Academy, passe à CoddyKit PRO. Le cours CUDA Academy comprend 4 leçons au total.

Qu'est-ce que j'apprendrai dans « Groupes coopératifs » ?

Une API moderne pour des portées de synchronisation flexibles. Tu pratiques CUDA Academy avec du code pratique que tu exécutes directement dans le navigateur, et un tuteur IA 24/7 répond à tes questions au fur et à mesure que tu avances dans la leçon.

Dois-je avoir de l'expérience pour commencer CUDA Academy ?

Aucune expérience préalable n'est requise. CUDA Academy sur CoddyKit est structuré pour les débutants jusqu'aux apprenants avancés, donc tu peux commencer ici ou depuis le début et avancer à ton rythme. Ceci est la leçon 4 sur 4.

Combien de temps prend la leçon « Groupes coopératifs » ?

La plupart des leçons CoddyKit prennent environ 5–10 minutes. Chacune est courte et interactive, tu progresses régulièrement et tu repiques exactement où tu t'es arrêté sur le web et l'app.

Peux-tu écrire et exécuter du code dans cette leçon CUDA Academy ?

Oui. Chaque leçon CUDA Academy inclut un éditeur de code intégré, tu écris et exécutes du vrai code directement dans ton navigateur et tu reçois des retours IA instantanés — aucune configuration locale requise.

Toutes les leçons de ce cours

  1. Warps, voies et masques
  2. __shfl_down_sync pour les réductions
  3. Fonctions de vote et de scrutin
  4. Groupes coopératifs
← Retour à CUDA Academy