0Pricing
CUDA Academy · บทเรียน

ใช้ __shfl_down_sync เพื่อลดรูป

ลดรูปภายในวาร์ปโดยไม่ใช้จุดกั้น

ใช้ __shfl_down_sync เพื่อลดรูป เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 2 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน

บางส่วนของบทเรียนนี้ยังไม่ได้รับการแปล และแสดงเป็นภาษาอังกฤษ

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

คำถามที่พบบ่อย

บทเรียน “ใช้ __shfl_down_sync เพื่อลดรูป” ฟรีหรือไม่

ใช่ — ข้อความเต็มของ “ใช้ __shfl_down_sync เพื่อลดรูป” ฟรีให้อ่านที่นี่บนเว็บ เพื่อปฏิบัติแบบโต้ตอบ (ตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7) และปลดล็อคส่วนที่เหลือของคอร์ส CUDA Academy ให้อัปเกรดเป็น CoddyKit PRO คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน

คุณจะเรียนรู้อะไรในบทเรียน “ใช้ __shfl_down_sync เพื่อลดรูป”

ลดรูปภายในวาร์ปโดยไม่ใช้จุดกั้น คุณปฏิบัติ CUDA Academy ด้วยโค้ดที่ใช้งานได้จริงที่คุณเรียกใช้โดยตรงในเบราว์เซอร์ และติวเตอร์ AI ตลอด 24/7 ตอบคำถามของคุณขณะที่คุณไปผ่านบทเรียน

คุณต้องมีประสบการณ์ก่อนที่จะเริ่มเรียน CUDA Academy หรือไม่

ไม่จำเป็นต้องมีประสบการณ์มาก่อน CUDA Academy บน CoddyKit ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 2 จากทั้งหมด 4 บทเรียน

บทเรียน “ใช้ __shfl_down_sync เพื่อลดรูป” ใช้เวลานานแค่ไหน

บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย

ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม

ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ

บทเรียนทั้งหมดในหลักสูตรนี้

  1. วาร์ป เลน และมาสก์
  2. ใช้ __shfl_down_sync เพื่อลดรูป
  3. ฟังก์ชัน Ballot และ Vote
  4. กลุ่มแบบร่วมมือ
← กลับไปที่ CUDA Academy