การอ่านแบบรวมกลุ่มเทียบกับแบบมีระยะก้าว

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

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

เธรด 0 เข้าถึงดัชนี 0 เธรด 1 เข้าถึงดัชนี 1 และต่อไปเรื่อย ๆ ที่อยู่ทั้ง 32 แห่งอยู่ในบรรทัดที่จัดแนวเดียวกัน ดังนั้น GPU จึงต้องใช้ธุรกรรมเดียว

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

เมื่อระยะก้าวเท่ากับ 2 เธรด 0 จะอ่าน 0 เธรด 1 จะอ่าน 2 และเธรด 2 จะอ่าน 4 ตอนนี้วาร์ปครอบคลุมหน่วยความจำกว้างขึ้นเป็นสองเท่า จึงต้องใช้ธุรกรรมมากขึ้น

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

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

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

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

เมื่อต้องประเมินรูปแบบ อย่านึกภาพเพียงเธรดเดียว ให้นึกถึงเลนทั้ง 32 เลนพร้อมกัน แล้วถามว่าที่อยู่ของพวกมันครอบคลุมกี่บรรทัดโดยรวม 🔍

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

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