รีจิสเตอร์และหน่วยความจำเฉพาะที่
พื้นที่จัดเก็บต่อเธรดที่เร็วที่สุดและข้อมูลที่ล้นออกมา
รีจิสเตอร์และหน่วยความจำเฉพาะที่ เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 1 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
หน่วยความจำที่เร็วที่สุดของคุณ
ทุกเธรดจะมีรีจิสเตอร์ส่วนตัวของตนเอง ซึ่งเป็นหน่วยจัดเก็บข้อมูลที่เร็วที่สุดบนชิป รีจิสเตอร์อยู่ติดกับหน่วยคำนวณ จึงแทบไม่มีค่าใช้จ่ายในการอ่าน
ตัวแปรทั่วไปกลายเป็นรีจิสเตอร์
เมื่อคุณเขียนตัวแปรภายในเครื่องทั่วไปในเคอร์เนล คอมไพเลอร์มักจะเก็บตัวแปรนั้นไว้ในรีจิสเตอร์โดยไม่ต้องใช้ไวยากรณ์พิเศษใด ๆ เลย 🙂
__global__ void k() {
int x = 5; // x lives in a register
float y = x * 2.0f;
}เป็นส่วนตัวสำหรับเธรดเดียว
ค่าของรีจิสเตอร์จะเป็นของเธรดเพียงหนึ่งเธรดเท่านั้น เธรดที่ 7 จะมองไม่เห็นรีจิสเตอร์ของเธรดที่ 8 นี่จึงทำให้งานของแต่ละเธรดแยกจากกันอย่างชัดเจน
รีจิสเตอร์เป็นกลุ่มทรัพยากรร่วม
มัลติโพรเซสเซอร์แบบสตรีมแต่ละตัวมีไฟล์รีจิสเตอร์ที่มีขนาดตายตัว เธรดที่ทำงานอยู่ทั้งหมดต้องแบ่งใช้พื้นที่นี้ ดังนั้นการใช้รีจิสเตอร์ต่อเธรดให้น้อยลงจึงทำให้มีเธรดทำงานพร้อมกันได้มากขึ้น
เมื่อรีจิสเตอร์ไม่พอ
หากเคอร์เนลต้องการรีจิสเตอร์มากกว่าที่มีอยู่ ค่าที่เกินมาจะถูกย้ายไปเก็บที่อื่น เหตุการณ์นี้เรียกว่ารีจิสเตอร์ล้น
รีจิสเตอร์ล้นไปที่ใด: หน่วยความจำภายในเครื่อง
ค่าที่ล้นจะไปอยู่ในหน่วยความจำภายในเครื่อง ซึ่งแม้ชื่อจะฟังดูเหมือนอยู่ใกล้ แต่จริง ๆ แล้วไม่ได้อยู่บนชิป แต่อยู่ใน DRAM นอกชิปที่ทำงานช้า
หน่วยความจำภายในเครื่องยังคงเป็นของแต่ละเธรด
หน่วยความจำภายในเครื่องเป็นพื้นที่ส่วนตัวของแต่ละเธรดเช่นเดียวกับรีจิสเตอร์ ความแตกต่างมีเพียงความเร็ว เพราะการเข้าถึงหน่วยความจำภายในเครื่องช้ากว่ามาก
อาร์เรย์ภายในเครื่องขนาดใหญ่ทำให้เกิดการล้น
การประกาศอาร์เรย์ขนาดใหญ่สำหรับแต่ละเธรดมักทำให้อาร์เรย์นั้นถูกย้ายไปไว้ในหน่วยความจำภายในเครื่อง เพราะมีรีจิสเตอร์ไม่เพียงพอที่จะเก็บทุกองค์ประกอบ
__global__ void k() {
float buf[64]; // likely lives in local memory
}รีจิสเตอร์ล้นทำให้ประสิทธิภาพลดลง
การล้นแต่ละครั้งเปลี่ยนการเข้าถึงรีจิสเตอร์ที่แทบไม่มีค่าใช้จ่ายให้กลายเป็นการเดินทางไปยัง DRAM ที่ช้า การลดแรงกดดันต่อรีจิสเตอร์เป็นหนึ่งในวิธีเพิ่มประสิทธิภาพที่ทำได้ง่ายที่สุด
ดูจำนวนรีจิสเตอร์ของคุณ
คุณสามารถขอให้คอมไพเลอร์รายงานจำนวนรีจิสเตอร์ต่อเธรดได้ ส่ง --ptxas-options=-v ให้ nvcc แล้วระบบจะแสดงการใช้งานของแต่ละเคอร์เนล
nvcc --ptxas-options=-v kernel.cuจำกัดจำนวนรีจิสเตอร์โดยตั้งใจ
ตัวระบุ __launch_bounds__ จะบอกเป็นนัยให้คอมไพเลอร์จำกัดจำนวนรีจิสเตอร์ โดยแลกความเร็วต่อเธรดเล็กน้อยกับการให้เธรดทำงานพร้อมกันได้มากขึ้น
__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }ตรวจสอบอย่างรวดเร็ว
คุณประกาศอาร์เรย์ขนาดใหญ่สำหรับแต่ละเธรด แล้วประสิทธิภาพลดลง สิ่งใดน่าจะเกิดขึ้น
ทบทวน: เร็วและเป็นส่วนตัว
คุณได้เรียนรู้ว่ารีจิสเตอร์เป็นพื้นที่จัดเก็บข้อมูลต่อเธรดที่เร็วที่สุด พื้นที่นี้มีจำกัด และข้อมูลที่เกินจะล้นไปยังหน่วยความจำภายในเครื่องที่ช้า ควรรักษาแรงกดดันต่อรีจิสเตอร์ให้ต่ำ 🎯
เรียนรู้ C++ ด้วย AI tutor — ฟรี
เขียนและเรียกใช้โค้ดจริงในเบราว์เซอร์ของคุณ รับความช่วยเหลือทันทีจาก AI tutor 24/7 และเรียนรู้ต่อจากที่คุณหยุดบนเว็บหรือในแอป
- คอร์ส
- 30
- บทเรียน
- 120
คำถามที่พบบ่อย
บทเรียน “รีจิสเตอร์และหน่วยความจำเฉพาะที่” ฟรีหรือไม่
ใช่ — ข้อความเต็มของ “รีจิสเตอร์และหน่วยความจำเฉพาะที่” ฟรีให้อ่านที่นี่บนเว็บ เพื่อปฏิบัติแบบโต้ตอบ (ตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7) และปลดล็อคส่วนที่เหลือของคอร์ส CUDA Academy ให้อัปเกรดเป็น CoddyKit PRO คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
คุณจะเรียนรู้อะไรในบทเรียน “รีจิสเตอร์และหน่วยความจำเฉพาะที่”
พื้นที่จัดเก็บต่อเธรดที่เร็วที่สุดและข้อมูลที่ล้นออกมา คุณปฏิบัติ CUDA Academy ด้วยโค้ดที่ใช้งานได้จริงที่คุณเรียกใช้โดยตรงในเบราว์เซอร์ และติวเตอร์ AI ตลอด 24/7 ตอบคำถามของคุณขณะที่คุณไปผ่านบทเรียน
คุณต้องมีประสบการณ์ก่อนที่จะเริ่มเรียน CUDA Academy หรือไม่
ไม่จำเป็นต้องมีประสบการณ์มาก่อน CUDA Academy บน CoddyKit ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 1 จากทั้งหมด 4 บทเรียน
บทเรียน “รีจิสเตอร์และหน่วยความจำเฉพาะที่” ใช้เวลานานแค่ไหน
บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย
ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม
ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ
บทเรียนทั้งหมดในหลักสูตรนี้
- รีจิสเตอร์และหน่วยความจำเฉพาะที่
- ข้อแลกเปลี่ยนของหน่วยความจำส่วนกลาง
- หน่วยความจำคงที่และแคชของมัน
- แบบจำลองทางความคิดของลำดับชั้น