ประกาศอาร์เรย์ __shared__

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

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

คุณประกาศใช้งานด้วยตัวระบุคุณสมบัติ __shared__ ภายในเคอร์เนล จากนั้นอาร์เรย์นี้จะถูกใช้ร่วมกันโดยทุกเธรดในบล็อก

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

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

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

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

รูปแบบที่พบบ่อยคือกำหนดหนึ่งช่องต่อเธรด โดยใช้ threadIdx.x เป็นดัชนี แต่ละเธรดจะโหลดองค์ประกอบของตนเข้าสู่หน่วยความจำที่ใช้ร่วมกันพร้อมกัน

หน่วยความจำที่ใช้ร่วมกันมีขนาดเล็ก โดยมักมีเพียง48 ถึง 100 KB ต่อ SM หากขอใช้ต่อบล็อกมากเกินไป จำนวนบล็อกที่ทำงานพร้อมกันได้จะลดลง

งานที่พบบ่อยที่สุดคือการเตรียมไทล์ของข้อมูลส่วนกลาง เพื่อให้บล็อกนำกลับมาใช้ซ้ำได้หลายครั้งโดยไม่ต้องกลับไปอ่าน DRAM ที่ช้า

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

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