ไปป์ไลน์แบบบัฟเฟอร์คู่
แบ่งข้อมูลเป็นส่วน ๆ เพื่อให้ GPU มีงานป้อนต่อเนื่อง
ไปป์ไลน์แบบบัฟเฟอร์คู่ เป็นบทเรียน CUDA Academy ฟรีบน CoddyKit นี่คือบทเรียนที่ 4 จากทั้งหมด 4 บทเรียน คุณสามารถอ่านบทเรียนทั้งหมดด้านล่างฟรี — จากนั้นลองปฏิบัติด้วยตัวคุณเองในเบราว์เซอร์พร้อมตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7 บทเรียนนี้เป็นส่วนหนึ่งของเส้นทางการเรียน CUDA Academy และความก้าวหน้าของคุณจะซิงค์ข้ามเว็บและแอป CoddyKit คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
ปัญหา GPU ว่างงาน
คัดลอกอาร์เรย์ขนาดใหญ่ จากนั้นเรียกใช้เคอร์เนล แล้วคัดลอกข้อมูลกลับ ระหว่างการคัดลอกแต่ละครั้ง GPU จะ ว่างงาน และระหว่างการคำนวณ เครื่องยนต์ถ่ายโอนก็จะว่างงาน
แบ่งงานออกเป็นส่วน ๆ
วิธีแก้เริ่มจากแบ่งอาร์เรย์ขนาดใหญ่ออกเป็น ส่วนย่อย ที่เล็กลง แต่ละส่วนสามารถคัดลอกและประมวลผลแยกกันได้ ทำให้มีโอกาสทำงานซ้อนทับกัน
สองบัฟเฟอร์ สองช่องทาง
การใช้บัฟเฟอร์สองชุดจะใช้ บัฟเฟอร์บนอุปกรณ์สองชุดและสตรีมสองรายการ ขณะที่สตรีม A คำนวณส่วนหนึ่ง สตรีม B จะคัดลอกส่วนถัดไปเข้าไป
ทำให้การคัดลอกซ้อนทับกับการคำนวณ
ความมหัศจรรย์อยู่ที่การทำงานพร้อมกัน คือการถ่ายโอนและ เคอร์เนลทำงานในเวลาเดียวกันกับข้อมูลคนละส่วน เวลาการถ่ายโอนจึงถูกซ่อนไว้เบื้องหลังการคำนวณที่เป็นประโยชน์
ต้องใช้หน่วยความจำแบบตรึง
วิธีนี้จะทำงานได้ก็ต่อเมื่อข้อมูลโฮสต์อยู่ในหน่วยความจำแบบ ตรึง บัฟเฟอร์ที่สลับหน้าได้จะบังคับให้ใช้การคัดลอกแบบบล็อก และกระบวนการทั้งหมดจะกลับไปทำงานแบบเรียงลำดับ
ลูปของกระบวนการต่อเนื่อง
คุณวนซ้ำไปตามส่วนย่อยต่าง ๆ และในแต่ละรอบจะสั่งคัดลอกเข้าแบบอะซิงโครนัส เรียกใช้ เคอร์เนล และคัดลอกออกแบบอะซิงโครนัส ทั้งหมดอยู่ในสตรีมเดียวกัน
cudaMemcpyAsync(d_in, h_in + off, csz, H2D, s);
proc<<<g, b, 0, s>>>(d_in, d_out);
cudaMemcpyAsync(h_out + off, d_out, csz, D2H, s);สลับสตรีม
กำหนดแต่ละส่วนย่อยให้กับสตรีมด้วยการ สลับสตรีม ส่วนที่มีหมายเลขคู่ไปยังสตรีม 0 และส่วนที่มีหมายเลขคี่ไปยังสตรีม 1 ทำให้ส่วนที่อยู่ติดกันทำงานซ้อนทับกันได้อย่างเหมาะสม
cudaStream_t s = streams[i % 2];เหตุใดสองสตรีมจึงเพียงพอ
สตรีมสองรายการก็ทำให้การคัดลอกซ้อนทับกับการคำนวณได้แล้ว บัฟเฟอร์ที่มากขึ้นจะเพิ่ม ความลึกของการทำงาน แต่ให้ผลตอบแทนลดลงและใช้หน่วยความจำเพิ่ม ดังนั้นควรเริ่มจากสองรายการ
รอให้งานทั้งหมดเสร็จตอนท้าย
หลังจากเข้าคิวส่วนย่อยทุกส่วนแล้ว ให้รอให้งานทั้งหมดเสร็จก่อนอ่านผลลัพธ์ การเรียกใช้ cudaDeviceSynchronize เพียงครั้งเดียวจะรอให้งานในทุกสตรีมเสร็จพร้อมกัน
cudaDeviceSynchronize();ความเร็วที่เพิ่มขึ้น
เมื่อซ่อนการคัดลอกไว้เบื้องหลังการคำนวณ เวลารวมจะลดลงจนใกล้เคียงกับต้นทุนของงานที่ นานกว่าเพียงงานเดียว ไม่ใช่ผลรวมของทั้งสองงาน นี่คือประโยชน์ที่แท้จริง 🚀
ได้ผลดีที่สุดเมื่อใด
การใช้บัฟเฟอร์สองชุดให้ผลดีที่สุดเมื่อเวลาถ่ายโอนและเวลาคำนวณ ใกล้เคียงกัน หากด้านใดด้านหนึ่งใช้เวลามากกว่ามาก ควรเริ่มลดเวลาของด้านนั้นก่อน
ตรวจสอบอย่างรวดเร็ว
ประโยชน์หลักของกระบวนการทำงานแบบใช้บัฟเฟอร์สองชุดคืออะไร
ทบทวน
แบ่งข้อมูลเป็นส่วนย่อย ใช้บัฟเฟอร์แบบตรึงและสตรีมสองรายการ แล้วทำให้การคัดลอกซ้อนทับกับการคำนวณ ซิงโครไนซ์ตอนท้าย แล้วดูเวลาถ่ายโอน หายไป 🎉
คำถามที่พบบ่อย
บทเรียน “ไปป์ไลน์แบบบัฟเฟอร์คู่” ฟรีหรือไม่
ใช่ — ข้อความเต็มของ “ไปป์ไลน์แบบบัฟเฟอร์คู่” ฟรีให้อ่านที่นี่บนเว็บ เพื่อปฏิบัติแบบโต้ตอบ (ตัวแก้ไขโค้ดในตัวและติวเตอร์ AI ตลอด 24/7) และปลดล็อคส่วนที่เหลือของคอร์ส CUDA Academy ให้อัปเกรดเป็น CoddyKit PRO คอร์ส CUDA Academy มีบทเรียนทั้งหมด 4 บทเรียน
คุณจะเรียนรู้อะไรในบทเรียน “ไปป์ไลน์แบบบัฟเฟอร์คู่”
แบ่งข้อมูลเป็นส่วน ๆ เพื่อให้ GPU มีงานป้อนต่อเนื่อง คุณปฏิบัติ CUDA Academy ด้วยโค้ดที่ใช้งานได้จริงที่คุณเรียกใช้โดยตรงในเบราว์เซอร์ และติวเตอร์ AI ตลอด 24/7 ตอบคำถามของคุณขณะที่คุณไปผ่านบทเรียน
คุณต้องมีประสบการณ์ก่อนที่จะเริ่มเรียน CUDA Academy หรือไม่
ไม่จำเป็นต้องมีประสบการณ์มาก่อน CUDA Academy บน CoddyKit ออกแบบมาสำหรับผู้เริ่มต้นไปจนถึงผู้เรียนขั้นสูง คุณสามารถเริ่มต้นที่นี่หรือเริ่มจากตัวแรกและเรียนด้วยความเร็วของคุณเอง นี่คือบทเรียนที่ 4 จากทั้งหมด 4 บทเรียน
บทเรียน “ไปป์ไลน์แบบบัฟเฟอร์คู่” ใช้เวลานานแค่ไหน
บทเรียน CoddyKit ส่วนใหญ่ใช้เวลาประมาณ 5–10 นาที แต่ละบทเรียนจึงสั้นและเป็นแบบโต้ตอบ คุณสามารถก้าวหน้าอย่างต่อเนื่องและกลับมาเรียนต่อจากตรงที่เพิ่งหยุดบนเว็บและแอปได้เลย
ฉันเขียนและรันโค้ดในบทเรียน CUDA Academy นี้ได้ไหม
ได้ บทเรียน CUDA Academy ทุกบทมีตัวแก้ไขโค้ดในตัว คุณจึงเขียนและรันโค้ดจริงได้เลยในเบราว์เซอร์ และได้รับข้อเสนอแนะจาก AI ในทันที — ไม่ต้องติดตั้งในเครื่องของคุณ
บทเรียนทั้งหมดในหลักสูตรนี้
- ทำไมหน่วยความจำแบบแบ่งหน้าได้จึงช้า
- หน่วยความจำตรึงด้วย cudaMallocHost
- cudaMemcpyAsync ในสตรีม
- ไปป์ไลน์แบบบัฟเฟอร์คู่