ซิงโครไนซ์ด้วย __syncthreads

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

จุดกั้นคือจุดในโค้ดที่เธรดทุกตัวต้องรอจนกว่าจะมาถึงครบทั้งหมด จากนั้นบล็อกจึงจะทำงานต่อไปพร้อมกัน

CUDA มีจุดกั้นระดับบล็อกชื่อ __syncthreads() ให้คุณใช้ เมื่อเรียกใช้ เธรดทุกตัวในบล็อกจะหยุดรอจนกว่าเธรดสุดท้ายจะมาถึง 🛑

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

หากลืมจุดกั้น เธรดหนึ่งอาจอ่านช่องข้อมูลก่อนที่เธรดข้างเคียงจะเขียนเสร็จ นี่คือสภาวะแข่งขัน และผลลัพธ์คือข้อมูลผิดที่เกิดขึ้นโดยไม่มีสัญญาณเตือน

จุดกั้นยังทำให้การเขียนลงหน่วยความจำที่ใช้ร่วมกันมองเห็นได้โดยเธรดทั้งหมดหลังจากผ่านจุดกั้น ดังนั้นเมื่อผ่าน __syncthreads แล้ว ทุกเธรดจะเห็นค่าที่เพิ่งเขียนใหม่

กฎที่เคร่งครัดคือ เธรดทุกตัวในบล็อกต้องมาถึง __syncthreads เดียวกัน หากบางเธรดข้ามจุดนี้ บล็อกอาจค้างตลอดไป

อย่าวางจุดกั้นไว้ภายในifที่มีเพียงบางเธรดเข้าไป เธรดที่เลือกเส้นทางอื่นจะมาไม่ถึง และบล็อกจะหยุดชะงักไม่สิ้นสุด ⚠️

ข้อจำกัดสำคัญประการหนึ่งคือ __syncthreads จะซิงโครไนซ์ได้เฉพาะบล็อกเดียวเท่านั้น ไม่สามารถประสานเธรดข้ามบล็อกต่าง ๆ ในกริดได้

ในอัลกอริทึมแบบไทล์ คุณมักต้องซิงก์สองครั้งต่อระยะการทำงาน: ครั้งหนึ่งหลังโหลดไทล์ และอีกครั้งหลังคำนวณเสร็จก่อนโหลดไทล์ถัดไป

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

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