CUDA Academy · บทเรียน

วาร์ป เลน และมาสก์

หน่วยเธรด 32 ตัวและมาสก์ที่กำลังทำงาน

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

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

เธรดทำงานเป็นเวิร์ป

GPU ไม่ได้จัดตารางเธรดทีละรายการ แต่จะจัดกลุ่มเป็น เวิร์ปละ 32 เธรดที่เคลื่อนที่ไปพร้อมกัน โดยทำงานตามคำสั่งเดียวกันแบบประสานจังหวะ

เหตุใด 32 จึงสำคัญ

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

แต่ละเธรดคือเลน

ภายในเวิร์ป แต่ละเธรดมีตำแหน่งตั้งแต่ 0 ถึง 31 เรียกว่า เลน หมายเลขเลนคือสิ่งที่คำสั่งพื้นฐานของเวิร์ปใช้ระบุว่าใครกำลังสื่อสารกับใคร

int lane = threadIdx.x % 32;

การค้นหารหัสเลน

คุณสามารถอ่านเลนได้โดยตรงจากรีจิสเตอร์พิเศษ แทนการคำนวณเอง รีจิสเตอร์ laneid จะเก็บค่า 0 ถึง 31 ของเธรดปัจจุบันไว้เสมอ

การทำงานแบบประสานจังหวะมีข้อควรระวัง

แต่ละเวิร์ปใช้ตัวนับคำสั่งร่วมกันหนึ่งตัว เมื่อคำสั่ง if ส่งเลนไปคนละเส้นทาง เวิร์ปจะ แยกเส้นทาง และทำงานแต่ละเส้นทางสลับกัน ส่งผลให้ความเร็วลดลง

หน้ากากที่กำลังทำงาน

ไม่ใช่ทุกเลนที่จะทำงานอยู่เสมอ หน้ากากขนาด 32 บิตจะระบุว่าเลนใดกำลังทำงาน โดยมีหนึ่งบิตต่อหนึ่งเลน และตั้งเป็น 1 เมื่อเลนนั้นมีส่วนร่วม

เหตุใดจึงมีหน้ากาก

หลังจากแยกเส้นทาง จะมีเพียงบางเลนที่ยังทำงานอยู่ คำสั่งพื้นฐานของเวิร์ปจึงต้องทราบอย่างแน่นอนว่าใครอยู่ในขณะนั้น คุณจึงต้องส่ง หน้ากากที่กำลังทำงานให้คำสั่งเหล่านั้นเพื่อให้ทำงานถูกต้อง

หน้ากากเต็มรูปแบบ

เมื่อแน่ใจว่าเลนทั้ง 32 เลนกำลังทำงาน หน้ากากจะเป็น 0xffffffff ซึ่งมีทุกบิตเป็น 1 หน้ากากเต็มรูปแบบนี้เป็นค่าที่ส่งให้คำสั่งบ่อยที่สุด

unsigned mask = 0xffffffff;

การสร้างหน้ากากอย่างปลอดภัย

ภายในแขนงคำสั่ง อย่าเดาหน้ากาก ให้เรียก activemask เพื่อบันทึกอย่างแม่นยำว่าเลนใดมาถึงจุดนี้ในขณะนั้น

unsigned mask = __activemask();

เลนสื่อสารกันได้โดยไม่ใช้หน่วยความจำ

ข้อดีสำคัญคือเลนในเวิร์ปเดียวกันสามารถแลกเปลี่ยนข้อมูลโดยตรงผ่านรีจิสเตอร์ได้ ไม่จำเป็นต้องใช้ หน่วยความจำใช้ร่วมหรือจุดกั้นสำหรับการแลกเปลี่ยนภายในเวิร์ป

การซิงก์หมายถึงระดับเวิร์ป

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

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

ลองทบทวนว่าเวิร์ปคืออะไร และมีขนาดเท่าใดบน GPU ของ NVIDIA

สรุปทบทวน

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

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

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

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

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

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

บทเรียน “วาร์ป เลน และมาสก์” ฟรีหรือไม่

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

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

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

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

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

บทเรียน “วาร์ป เลน และมาสก์” ใช้เวลานานแค่ไหน

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

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

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

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

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