0Pricing
CUDA Academy · Lección

__shfl_down_sync para reducciones

Realice reducciones de warps sin barreras.

__shfl_down_sync para reducciones es una lección gratuita de CUDA Academy en CoddyKit. Esta es la lección 2 de 4. Puedes leer la lección completa abajo gratuitamente — luego la practicas en el navegador con un editor de código integrado y un tutor de IA 24/7. Forma parte de la ruta de aprendizaje de CUDA Academy, y tu progreso se sincroniza en la web y la app de CoddyKit. El curso de CUDA Academy incluye 4 lecciones en total.

Partes de esta lección aún no han sido traducidas y se muestran en inglés.

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. ✨

Preguntas frecuentes

¿La lección «__shfl_down_sync para reducciones» es gratis?

Sí — el texto completo de «__shfl_down_sync para reducciones» es gratis para leer aquí en la web. Para practicarla de forma interactiva (editor de código integrado y tutor de IA 24/7) y desbloquear el resto del curso de CUDA Academy, actualiza a CoddyKit PRO. El curso de CUDA Academy incluye 4 lecciones en total.

¿Qué aprenderé en «__shfl_down_sync para reducciones»?

Realice reducciones de warps sin barreras. Practicas CUDA Academy con código real que ejecutas directamente en el navegador, y un tutor de IA 24/7 responde tus preguntas mientras trabajas en la lección.

¿Necesito experiencia previa para empezar CUDA Academy?

No se requiere experiencia previa. CUDA Academy en CoddyKit está estructurado para principiantes hasta estudiantes avanzados, así que puedes empezar aquí o desde el inicio y avanzar a tu ritmo. Esta es la lección 2 de 4.

¿Cuánto tiempo toma la lección «__shfl_down_sync para reducciones»?

La mayoría de las lecciones de CoddyKit toman alrededor de 5–10 minutos. Cada una es compacta e interactiva, así que avanzas constantemente y retomas exactamente por donde dejaste en la web y la app.

¿Puedo escribir y ejecutar código en esta lección de CUDA Academy?

Sí. Cada lección de CUDA Academy incluye un editor de código integrado, así que escribes y ejecutas código real directamente en tu navegador y obtienes retroalimentación instantánea de IA — sin configuración local necesaria.

Todas las lecciones de este curso

  1. Warps, lanes y máscaras
  2. __shfl_down_sync para reducciones
  3. Funciones de votación y ballot
  4. Cooperative Groups
← Volver a CUDA Academy