การระบุแอดเดรสตามลำดับ
ระยะก้าวในหน่วยความจำร่วมที่ไม่ก่อความขัดแย้ง
การระบุแอดเดรสตามลำดับ เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 3 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
หน่วยความจำใช้ร่วมกันแบ่งเป็นแบงก์
หน่วยความจำใช้ร่วมกันแบ่งออกเป็น 32 แบงก์ แบงก์ละหนึ่งเลนของวาร์ป เมื่อเธรดทั้ง 32 เข้าถึงแบงก์ที่แตกต่างกัน 32 แบงก์ การอ่านทั้งหมดจะเกิดขึ้นภายในรอบการทำงานเดียวที่รวดเร็ว
ความขัดแย้งของแบงก์ทำให้ช้าลง
หากเธรดสองเธรดในวาร์ปเข้าถึงแบงก์เดียวกัน นั่นคือ ความขัดแย้งของแบงก์ ฮาร์ดแวร์จะทำให้การเข้าถึงเหล่านั้นทำงานทีละรายการ จึงต้องใช้รอบการทำงานเพิ่มขึ้น
การกำหนดแอดเดรสแบบสลับ
การลดค่าก่อนหน้านี้ใช้การกำหนดแอดเดรสแบบ สลับ โดยเริ่มจากระยะก้าวที่เล็กแล้วเพิ่มเป็นสองเท่าในแต่ละขั้น ทำให้คู่ข้อมูลอยู่ใกล้กันในหน่วยความจำใช้ร่วมกัน
int index = 2 * s * tid;
data[index] += data[index + s];เหตุใดการสลับจึงทำให้เกิดความขัดแย้ง
เมื่อระยะก้าวมีขนาดเล็กและเพิ่มเป็นสองเท่า เลนหลายเลนในวาร์ปจะถูกจับคู่กับ แบงก์ เดียวกัน การเข้าถึงเหล่านั้นจึงไม่สามารถเกิดขึ้นภายในรอบเดียวได้อีกต่อไป
สลับลำดับของระยะก้าว
การกำหนดแอดเดรสแบบลำดับจะเริ่มด้วยระยะก้าวที่ ใหญ่ แล้วลดลงครึ่งหนึ่งในแต่ละขั้น ซึ่งตรงข้ามกับการสลับ การเปลี่ยนแปลงเพียงอย่างเดียวนี้จะกำจัดความขัดแย้งได้
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s)
data[tid] += data[tid + s];
__syncthreads();
}ระยะก้าวใหญ่ แบงก์สะอาด
ระยะก้าวขนาดใหญ่ทำให้แอดเดรสของคู่ข้อมูลอยู่ห่างกันมาก ดังนั้นแต่ละเลนจึงตกลงบน แบงก์ของตนเอง วาร์ปจึงอ่านได้โดยไม่มีความขัดแย้งภายในรอบเดียว
เงื่อนไข tid < s
ในแต่ละขั้นจะมีเพียงเธรดครึ่งล่างเท่านั้นที่ทำงาน ซึ่งเขียนเป็น tid < s วิธีนี้ทำให้เธรดที่ทำงานอยู่เรียงต่อเนื่องกัน และวาร์ปก็ไม่เกิดการแตกแขนงเช่นกัน
ได้ประโยชน์สองอย่างพร้อมกัน
การกำหนดแอดเดรสแบบลำดับแก้ไข ความขัดแย้งของแบงก์ และหลีกเลี่ยงการแตกแขนงของวาร์ปภายในเคอร์เนลเดียวกัน การเปลี่ยนรูปแบบเพียงครั้งเดียวแก้ปัญหาด้านประสิทธิภาพได้สองอย่าง
ยังต้องซิงโครไนซ์ทุกขั้น
คุณยังคงต้องเรียกใช้ __syncthreads หลังแต่ละขั้น เธรดต้องมองเห็นการเขียนของระดับก่อนหน้า ก่อนจะอ่านข้อมูลสำหรับระดับถัดไป
ผลลัพธ์ไปอยู่ที่ดัชนี 0
เมื่อระยะก้าวลดลงครึ่งหนึ่งจนเข้าใกล้ศูนย์ ผลรวมย่อยทั้งหมดจะถูกรวมเข้าที่ data[0] จากนั้นเธรด 0 จะเขียนผลลัพธ์ของบล็อกนั้นกลับไปยังหน่วยความจำส่วนกลาง
การปรับให้เหมาะสมแบบคลาสสิก
รูปแบบนี้มาจากคู่มือการลดค่าชื่อดังของ NVIDIA โดยตรง การกำหนดแอดเดรสแบบลำดับเป็นขั้นพื้นฐานสู่เคอร์เนลที่ ไม่มีความขัดแย้ง
ตรวจสอบความเข้าใจอย่างรวดเร็ว
ลองคิดดูว่าเหตุใดระยะก้าวขนาดใหญ่ที่ลดลงครึ่งหนึ่งจึงดีกว่าระยะก้าวขนาดเล็กที่เพิ่มเป็นสองเท่า
สรุป
คุณเปลี่ยนจากการกำหนดแอดเดรสแบบสลับเป็น การกำหนดแอดเดรสแบบลำดับ โดยเริ่มด้วยระยะก้าวขนาดใหญ่แล้วลดลงครึ่งหนึ่ง จึงกำจัดทั้งความขัดแย้งของแบงก์และการแตกแขนงได้พร้อมกัน ต่อไปคือผลรวมจากหลายบล็อก ✨
เรียนรู้ 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 ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 3 จากทั้งหมด 4 บทเรียน
บทเรียน “การระบุแอดเดรสตามลำดับ” ใช้เวลานานแค่ไหน
บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย
ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม
ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ
บทเรียนทั้งหมดในหลักสูตรนี้
- แนวคิดต้นไม้ลดรูป
- กำจัดการแยกทางของวาร์ป
- การระบุแอดเดรสตามลำดับ
- การลดรูปขั้นสุดท้ายหลายบล็อก