0Pricing
CUDA Academy · Ders

Azaltmalar için __shfl_down_sync

Engeller olmadan warp azaltmaları.

Azaltmalar için __shfl_down_sync, CoddyKit'te ücretsiz bir CUDA Academy dersidir. Bu, 4 dersinin 2. dersidir. Aşağıdan dersin tamamını ücretsiz okuyabilir, sonra tarayıcıda yerleşik kod editörü ve 7/24 yapay zeka koçu ile uygulamalı olarak pratik yapabilirsin. Bu, CUDA Academy öğrenme yolunun bir parçasıdır ve ilerlemeniz web ve CoddyKit uygulaması arasında senkronize olur. CUDA Academy kursu toplamda 4 dersten oluşur.

Bu dersin bazı bölümleri henüz çevrilmemiş olup İngilizce olarak gösterilmektedir.

Shuffle Moves Registers

A shuffle lets one lane read another lane's register value directly. Data hops between threads with no shared memory and no barrier in between.

Meet shfl_down_sync

The workhorse for reductions is shfl_down_sync. Each lane grabs a value from a lane a fixed number of positions higher in the warp.

float v = __shfl_down_sync(mask, val, offset);

Reading the Arguments

The call takes an active mask, the value to share, and an offset. Lane i receives the value held by lane i plus offset.

Add What You Receive

A warp sum just adds the shuffled value back in. Lane i takes its neighbor's number and folds it into its own running total.

val += __shfl_down_sync(mask, val, offset);

Halve the Offset

Like a tree reduction, the offset starts at 16 and halves each step: 16, 8, 4, 2, 1. After five steps the whole warp is summed.

for (int o = 16; o > 0; o >>= 1)
  val += __shfl_down_sync(mask, val, o);

Why Five Steps

A warp has 32 lanes, and 2 to the fifth power is 32. So five halving steps are exactly enough to combine all 32 values into one.

Answer Lands in Lane 0

After the loop, the complete warp total sits in lane 0. The other lanes hold partial junk, so only lane 0 should write the result out.

No Barriers Needed

Because exchange happens through registers in lockstep, you skip syncthreads entirely. The sync suffix already coordinates the masked lanes.

Faster Than Shared

A warp shuffle reduction beats the shared-memory version: fewer instructions, no bank conflicts, and no barrier stalls. It is the fast path for the last 32.

Pass the Right Mask

If some lanes already exited, a full mask is wrong. Capture the live set with activemask so each shuffle only touches present lanes.

Two-Level Reductions

Big reductions often reduce each warp with shuffles, then combine warp results in shared memory. Shuffles handle the cheap inner level beautifully.

Quick Check

Think about which lane holds the answer when the loop finishes.

Recap

You summed a warp with shfl_down_sync: halve the offset five times, and lane 0 holds the total, no barriers needed. Next: ballot and vote. ✨

Sıkça Sorulan Sorular

“Azaltmalar için __shfl_down_sync” dersi ücretsiz mi?

Evet — “Azaltmalar için __shfl_down_sync” dersin tüm metni burada web'de ücretsiz olarak okunabilir. Etkileşimli olarak pratik yapmak (yerleşik kod editörü ve 7/24 yapay zeka koçu) ve CUDA Academy kursunun geri kalanını açmak için CoddyKit PRO'ya yükselt. CUDA Academy kursu toplamda 4 dersten oluşur.

“Azaltmalar için __shfl_down_sync” dersinde ne öğreneceğim?

Engeller olmadan warp azaltmaları. CUDA Academy ile uygulamalı kodu tarayıcıda doğrudan çalıştırarak pratik yaparsın ve 7/24 yapay zeka koçu dersi çalışırken sorularını yanıtlar.

CUDA Academy öğrenmeye başlamak için deneyim gerekli mi?

Önceden deneyim gerekmez. CoddyKit'te CUDA Academy, başlangıçtan ileri seviyeye kadar yapılandırıldığı için buradan başlayabilir veya başından başlayıp kendi hızında ilerleme yapabilirsin. Bu, 4 dersinin 2. dersidir.

“Azaltmalar için __shfl_down_sync” dersi ne kadar sürer?

Çoğu CoddyKit dersi yaklaşık 5–10 dakika sürer. Her biri kısa ve etkileşimli olduğu için sabit ilerleme yaparsın ve web ile uygulama arasında tam olarak bıraktığın yerden devam edebilirsin.

Bu CUDA Academy dersinde kod yazıp çalıştırabilir miyim?

Evet. Her CUDA Academy dersi yerleşik bir kod editörü içerir, bu sayede tarayıcıda gerçek kod yazıp çalıştırabilir ve anlık yapay zeka geri bildirimi alırsın — yerel kurulum gerekli değildir.

Bu kursun tüm dersleri

  1. Warp'lar, Şeritler ve Maskeler
  2. Azaltmalar için __shfl_down_sync
  3. Oylama ve Oy İşlevleri
  4. İşbirlikçi Gruplar
← CUDA Academy Sayfasına Dön