CUDA Academy · บทเรียน

กลุ่มแบบร่วมมือ

API สมัยใหม่สำหรับขอบเขตการซิงโครไนซ์ที่ยืดหยุ่น

บทเรียน 4 จาก 413 ขั้นตอน

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

API การประสานงานที่ชัดเจนยิ่งขึ้น

การใช้ mask ดิบและอินทรินสิกทำได้ แต่ค่อนข้างยุ่งยาก กลุ่มทำงานร่วมกัน จะห่อหุ้มเธรดไว้เป็นออบเจ็กต์ที่คุณตั้งชื่อ กำหนดขนาด และประสานงานได้อย่างชัดเจน

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

ดึงทั้งบล็อก

เริ่มต้นด้วยการรับตัวอ้างอิงไปยังบล็อกเธรดของคุณ this_thread_block จะคืนค่ากลุ่มที่คุณสามารถสั่งให้ประสานงานได้เหมือน syncthreads แต่มีรูปแบบเป็นออบเจ็กต์

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

ประสานงานผ่านกลุ่ม

การเรียกใช้ sync กับบล็อกคือจุดกั้นรูปแบบใหม่ โดยทำงานเหมือน syncthreads ทุกประการ แต่สื่อความหมายได้ชัดเจนว่าเป็นเมธอดของกลุ่มที่คุณต้องการ

block.sync();

แบ่งเป็นส่วนย่อยขนาดวาร์ป

คุณสามารถแบ่งบล็อกออกเป็น ส่วนย่อย ที่มีขนาดคงที่ได้ ส่วนย่อยขนาด 32 เลนจะให้กลุ่มขนาดเท่าวาร์ป พร้อมเมธอดที่ชัดเจนแทนการใช้ mask ดิบสำหรับการสลับข้อมูล

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

เมธอดมาแทนที่ mask

กลุ่มแบบแบ่งส่วนมี shfl_down และเมธอดอื่น ๆ โดยไม่ต้องรับอาร์กิวเมนต์ mask เพราะกลุ่มรู้จักสมาชิกของตนเองอยู่แล้ว API จึงสั้นและปลอดภัย

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 รองรับความสามารถนี้ การเปิดใช้งานแบบปกติจะไม่อนุญาตให้ทำเช่นนั้น

ทำไมจึงควรใช้

กลุ่มทำงานร่วมกันทำให้โค้ดระดับวาร์ป อ่านเข้าใจได้ง่าย และย้ายไปใช้ได้กับหลายระบบ เพราะไม่ต้องจัดการ mask เอง มีขอบเขตที่ชัดเจน และมีการดำเนินการรวมที่นำกลับมาใช้ใหม่ได้ตรงกับความหมายที่ต้องการ

ตรวจสอบอย่างรวดเร็ว

ลองนึกดูว่าคุณสร้างกลุ่มทำงานร่วมกันขนาดเท่าวาร์ปจากบล็อกได้อย่างไร

สรุปทบทวน

กลุ่มทำงานร่วมกันเปลี่ยน mask ให้เป็นออบเจ็กต์ที่มีชื่อ คุณสามารถแบ่งบล็อกเป็นส่วนย่อย เรียกใช้ reduce หรือแม้แต่ประสานงานทั้งกริดได้ ตอนนี้คุณควบคุม CUDA ระดับวาร์ปได้แล้ว ✨

เริ่มต้นได้ฟรี

เรียนรู้ C++ ด้วย AI tutor — ฟรี

เขียนและเรียกใช้โค้ดจริงในเบราว์เซอร์ของคุณ รับความช่วยเหลือทันทีจาก AI tutor 24/7 และเรียนรู้ต่อจากที่คุณหยุดบนเว็บหรือในแอป

คอร์ส
30
บทเรียน
120

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

บทเรียน “กลุ่มแบบร่วมมือ” ฟรีหรือไม่

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

คุณจะเรียนรู้อะไรในบทเรียน “กลุ่มแบบร่วมมือ”

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

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

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

บทเรียน “กลุ่มแบบร่วมมือ” ใช้เวลานานแค่ไหน

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

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

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

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

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