CUDA Academy · บทเรียน

การระบุแอดเดรสตามลำดับ

ระยะก้าวในหน่วยความจำร่วมที่ไม่ก่อความขัดแย้ง

บทเรียน 3 จาก 413 ขั้นตอน

การระบุแอดเดรสตามลำดับ เป็นบทเรียน 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 ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ

บทเรียนทั้งหมดในหลักสูตรนี้

  1. แนวคิดต้นไม้ลดรูป
  2. กำจัดการแยกทางของวาร์ป
  3. การระบุแอดเดรสตามลำดับ
  4. การลดรูปขั้นสุดท้ายหลายบล็อก
← กลับไปที่ CUDA Academy