แรงกดดันของรีจิสเตอร์และการสปิล

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

เธรดที่ทำงานอยู่ทุกเธรดใช้ ไฟล์รีจิสเตอร์ ร่วมกันหนึ่งชุดต่อ SM ยิ่งแต่ละเธรดต้องใช้รีจิสเตอร์มากเท่าใด ก็ยิ่งมีเธรดอยู่พร้อมกันได้น้อยลงเท่านั้น

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

เมื่อเธรดต้องการรีจิสเตอร์มากกว่าที่มี คอมไพเลอร์จะย้ายค่าเพิ่มเติมออกไป การล้นนี้เรียกว่า การล้นรีจิสเตอร์ ไปยังหน่วยความจำที่ช้ากว่า

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

ขอให้คอมไพเลอร์รายงานการใช้งาน การเพิ่ม --ptxas-options=-v ให้กับ nvcc จะแสดงจำนวนรีจิสเตอร์ต่อเธรดและจำนวนไบต์ที่ล้นสำหรับเคอร์เนลแต่ละตัว

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

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

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

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

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

เมื่อเคอร์เนลใช้รีจิสเตอร์จนหมด ค่าเหล่านั้นจะถูกส่งไปที่ใด

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