ในระบบนิเวศการคำนวณด้วย GPU ในปัจจุบัน NVIDIA CCCL (CUDA Core Compute Libraries) ถือเป็นรากฐานระดับล่างที่ไม่สามารถมองข้ามได้

เป้าหมายของ CCCL คือการจัดเตรียมชุดไลบรารีคอมโพเนนต์พื้นฐานสำหรับนักพัฒนา CUDA C++ โดยมีจุดมุ่งหมายเพื่อลดความซับซ้อนในการเขียนโค้ดที่มีประสิทธิภาพและปลอดภัย การผสานรวมไลบรารีเหล่านี้สามารถลดความซับซ้อนในการพัฒนาอย่างมีนัยสำคัญ และใช้ประโยชน์จากประสิทธิภาพอันทรงพลังของ CUDA C++ ได้อย่างเต็มที่

ตั้งแต่เฟรมเวิร์กการอนุมาน SGLang, vLLM ไปจนถึงเฟรมเวิร์กการฝึก PyTorch, คอมไพเลอร์การเรียนรู้เชิงลึก TVM และไลบรารีการคำนวณทางวิทยาศาสตร์ RAPIDS โครงการ CUDA C++ กระแสหลักเกือบทั้งหมดล้วนพึ่งพา Thrust, CUB และ libcudacxx ที่ CCCL จัดเตรียมไว้ ไม่ทางตรงก็ทางอ้อม

สถานะของมันในระบบนิเวศ CUDA ก็เหมือนกับไลบรารีมาตรฐาน C++ ที่มีต่อ C++ เอง — เป็นโครงสร้างพื้นฐานทั่วไปสำหรับการคำนวณแบบขนาน

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

สารบัญบทความ

• หนึ่ง: ตำแหน่งโครงการและแรงจูงใจในการออกแบบ
• สอง: สถาปัตยกรรมนามธรรมสามชั้น
  • 2.1 libcudacxx — ชั้นพื้นฐาน: การนำไลบรารีมาตรฐาน C++ ไปใช้บน CUDA
  • 2.2 CUB — ชั้นกลาง: ไพรมิทีฟประสิทธิภาพสูงเฉพาะสำหรับ CUDA
  • 2.3 Thrust — ชั้นสูง: ไลบรารีอัลกอริทึมแบบขนานข้ามแบ็กเอนด์
• สาม: ระบบการสร้างแบบรวม
• สี่: การจัดการเนมสเปซและเวอร์ชัน
  • 4.1 กลไกการห่อหุ้มเนมสเปซ
  • 4.2 การควบคุมเวอร์ชันแบบรวม
  • 4.3 การจัดการเวอร์ชัน ABI
• ห้า: สรุปสถาปัตยกรรมโดยรวม

หนึ่ง: ตำแหน่งโครงการและแรงจูงใจในการออกแบบ

CCCL เป็นโครงการของ NVIDIA ที่รวมไลบรารีหลักสามไลบรารีของ CUDA C++ ที่พัฒนาอย่างอิสระ — Thrust, CUB, libcudacxx — เข้าสู่ที่เก็บซอร์สโค้ดเดียว

พันธกิจของโครงการถูกระบุไว้อย่างชัดเจนในเอกสารโครงการว่า “ทำให้ CUDA น่าพึงพอใจยิ่งขึ้น” โดยมีเป้าหมายเพื่อจัดเตรียมเครื่องมือพื้นฐานทั่วไปที่มีประสิทธิภาพสูงสุดสำหรับนักพัฒนา CUDA C++ ในบทบาทที่คล้ายคลึงกับไลบรารีมาตรฐาน C++ ที่มีต่อ C++ มาตรฐาน: ชุดเครื่องมือพื้นฐานทั่วไปที่มีประสิทธิภาพสูงสุด

ก่อนที่จะรวมเข้าด้วยกัน ไลบรารีทั้งสามได้รับการบำรุงรักษาอย่างอิสระ ส่งผลให้เกิดปัญหาการพัฒนาซ้ำซ้อน เวอร์ชันไม่สอดคล้องกัน และความซับซ้อนในการผสานรวมสำหรับผู้ใช้ CCCL 2.2.0 เป็นเวอร์ชันแรกที่เผยแพร่จากที่เก็บซอร์สโค้ดแบบรวม หลังจากนั้นคอมโพเนนต์ทั้งหมดจะใช้หมายเลขเวอร์ชันความหมาย (semantic version) ร่วมกัน

สอง: สถาปัตยกรรมนามธรรมสามชั้น

การออกแบบหลักของ CCCL แสดงให้เห็นนามธรรมสามชั้นที่ชัดเจน จากระดับล่างไปสูงตามลำดับ:

2.1 libcudacxx — ชั้นพื้นฐาน: การนำไลบรารีมาตรฐาน C++ ไปใช้บน CUDA

libcudacxx เป็นรากฐานของ CCCL ทั้งหมด จัดเตรียมการนำไลบรารีมาตรฐาน C++ ไปใช้ที่สามารถทำงานในโค้ดโฮสต์ (host) และอุปกรณ์ (device) พร้อมกันได้ รวมถึงนามธรรมเฉพาะฮาร์ดแวร์ CUDA (เช่น ไพรมิทีฟซิงโครไนซ์ การควบคุมแคช การดำเนินการอะตอมมิก ฯลฯ)

ไฟล์เฮดเดอร์ของมันถูกจัดระเบียบภายใต้ libcudacxx/include/cuda/ ครอบคลุมโมดูลฟังก์ชันที่หลากหลาย:

| ไดเรกทอรี | ฟังก์ชัน | ไฟล์เฮดเดอร์สาธารณะหลัก |
| :— | :— | :— |
| __atomic/ | การดำเนินการอะตอมมิก | cuda/atomic |
| __barrier/ | ซิงโครไนซ์แบบแบร์ริเออร์ | cuda/barrier |
| __memory_resource/ | การจัดการทรัพยากรหน่วยความจำ | cuda/memory_resource |
| __ptx/ | การห่อหุ้มคำสั่ง PTX | cuda/ptx |
| __stream/ | นามธรรมสตรีม CUDA | cuda/stream |
| __mdspan/ | มุมมองอาร์เรย์หลายมิติ | cuda/mdspan |
| __warp/ | ไพรมิทีฟระดับ Warp | cuda/warp |
| __launch/ | นามธรรมการลอจิกเคอร์เนล | cuda/launch |
| __execution/ | นโยบายการดำเนินการและการควบคุมความแน่นอน | — |
| std/ | ชั้นความเข้ากันได้กับไลบรารีมาตรฐาน C++ | cuda/std/* |

2.1.1 โมเดลขอบเขตของการดำเนินการอะตอมมิก

ยกตัวอย่างการดำเนินการอะตอมมิก ไฟล์เฮดเดอร์สาธารณะ cuda/atomic ใช้กลยุทธ์การรวมสองครั้งเพื่อให้เกิดความเข้ากันได้:

cpp

include // การนำ CUDA ไปใช้เฉพาะ (พร้อม thread_scope)

include // อินเทอร์เฟซความเข้ากันได้กับไลบรารีมาตรฐาน C++

คลาสหลัก cuda::atomic<T, Sco> ถูกกำหนดไว้ใน libcudacxx/include/cuda/__atomic/atomic.h มันสืบทอดมาจาก cuda::std::__atomic_impl และขยายการดำเนินการอะตอมมิกเฉพาะของ CUDA fetch_max และ fetch_min:

cpp
template
struct atomic : public ::cuda::std::__atomic_impl<_Tp, _Sco>
{
// อินเทอร์เฟซ atomic มาตรฐานสืบทอดผ่านคลาสฐาน
_CCCL_API constexpr atomic(_Tp __d) noexcept
: ::cuda::std::__atomic_impl<_Tp, _Sco>(__d) {}

// ส่วนขยาย CUDA: atomic fetch_max / fetch_min
_CCCL_API inline _Tp fetch_max(const _Tp& __op, memory_order __m = memory_order_seq_cst) noexcept
  {
    return ::cuda::std::__atomic_fetch_max_dispatch(
      &this->__a, __op, __m, ::cuda::std::__scope_to_tag<_Sco>{});
  }
// ... fetch_min ถูกนำมาใช้ในทำนองเดียวกัน

};

cuda::atomic_ref<T, Sco> คลาสที่สอดคล้องกัน จัดเตรียมการอ้างอิงอะตอมมิกไปยังตัวแปรที่มีอยู่ ซึ่งถูกกำหนดไว้ที่บรรทัด L79-L111 ในไฟล์เดียวกัน

แบร์ริเออร์เธรด (atomic_thread_fence) การนำไปใช้แสดงให้เห็นว่า libcudacxx จับคู่อินเทอร์เฟซนามธรรมกับไพรมิทีฟ CUDA ระดับล่างที่แตกต่างกันตามขอบเขตอย่างไร:

cpp
_CCCL_API inline void
atomic_thread_fence(memory_order __m, thread_scope _Scope = thread_scope::thread_scope_system)
{
NV_DISPATCH_TARGET(
NV_IS_DEVICE,
(switch (_Scope) {
case thread_scope::thread_scope_system:
::cuda::std::__atomic_thread_fence_cuda((int) __m, __thread_scope_system_tag{});
break;
case thread_scope::thread_scope_device:
::cuda::std::__atomic_thread_fence_cuda((int) __m, __thread_scope_device_tag{});
break;
case thread_scope::thread_scope_block:
::cuda::std::__atomic_thread_fence_cuda((int) __m, __thread_scope_block_tag{});
break;
case thread_scope::thread_scope_thread:
break; // ขอบเขตระดับเธรดไม่ต้องการ fence
}),
NV_IS_HOST,
(::cuda::std::atomic_thread_fence(__m);)) // ฝั่งโฮสต์ย้อนกลับไปใช้การนำมาตรฐานไปใช้
}

ในที่นี้ NV_DISPATCH_TARGET แมโครจะกระจายเส้นทางโค้ดอุปกรณ์หรือโฮสต์ในเวลาคอมไพล์ตามแพลตฟอร์มเป้าหมาย เพื่อขจัดค่าใช้จ่ายสาขาในรันไทม์

2.1.2 ซิงโครไนซ์แบบแบร์ริเออร์ (cuda::barrier)

cuda::barrier<Sco, CompletionF> ถูกกำหนดไว้ใน libcudacxx/include/cuda/__barrier/barrier.h มันยังควบคุมขอบเขตการซิงโครไนซ์อย่างละเอียดผ่านพารามิเตอร์เทมเพลต thread_scope:

cpp
template class barrier : public ::cuda::std::__barrier_base<_CompletionF, _Sco>
{
public:
_CCCL_API constexpr barrier(::cuda::std::ptrdiff_t __expected, _CompletionF __completion = _CompletionF())
: ::cuda::std::__barrier_base<_CompletionF, _Sco>(__expected, __completion) {}

// ฟังก์ชันเฟรนด์สำหรับการเริ่มต้นแบบอินเพลซ (placement new) ของ barrier ในหน่วยความจำ shared
_CCCL_API inline friend void init(barrier* __b, ::cuda::std::ptrdiff_t __expected) {
new (__b) barrier(__expected);
}
};

2.1.3 การห่อหุ้มคำสั่ง PTX

ไฟล์เฮดเดอร์เช่น libcudacxx/include/cuda/__ptx/instructions/get_sreg.h จัดเตรียมการห่อหุ้ม C++ สำหรับรีจิสเตอร์พิเศษ PTX (เช่น cuda::ptx::get_sreg_laneid()) การนำไปใช้ระดับล่างของไลบรารีระดับบนเช่น CUB จะใช้การห่อหุ้มเหล่านี้โดยตรง แทนที่จะเขียนแอสเซมบลีอินไลน์ดิบ

เมื่อนักพัฒนาใช้งาน โดยการระบุขอบเขต (เช่น cuda::thread_scope_device) และลำดับหน่วยความจำ (เช่น cuda::memory_order_relaxed) การดำเนินการอะตอมมิกจะถูกคอมไพล์เป็นคำสั่งอะตอมมิกฮาร์ดแวร์ GPU ที่สอดคล้องกันโดยตรง:

cpp
cuda::atomic_ref<int, cuda::thread_scope_device> atomic_result(result.front());
atomic_result.fetch_add(sum, cuda::memory_order_relaxed);

2.2 CUB — ชั้นกลาง: ไพรมิทีฟประสิทธิภาพสูงเฉพาะสำหรับ CUDA

CUB เป็นไลบรารีระดับล่างสำหรับนักพัฒนาเคอร์เนล CUDA จัดเตรียมไพรมิทีฟอัลกอริทึมแบบขนานประสิทธิภาพสูงระดับ บล็อก, วาร์ป และ อุปกรณ์

2.2.1 ไพรมิทีฟระดับบล็อก: ยกตัวอย่าง BlockReduce

ไพรมิทีฟระดับบล็อกเป็นคุณลักษณะหลักของ CUB cub/cub/block/block_reduce.cuh จัดเตรียมสี่รูปแบบอัลกอริทึมผ่านการแจงนับ BlockReduceAlgorithm (L43-L153):

3.2.1 BlockReduce: การวิเคราะห์รูปแบบอัลกอริทึมและการนำไปใช้

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

| รูปแบบอัลกอริทึม | ลักษณะหลัก | ขั้นตอนการดำเนินการ |
| :— | :— | :— |
| BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY | อัลกอริทึม raking ที่มีประสิทธิภาพ รองรับเฉพาะการดำเนินการที่สลับที่ได้ (เช่น การบวก การหาค่าสูงสุด) | ① รีดักชันรีจิสเตอร์ → ② รีดักชัน raking หน่วยความจำร่วม → ③ รีดักชัน Kogge-Stone ภายในวาร์ป |
| BLOCK_REDUCE_RAKING | อัลกอริทึม raking ทั่วไปที่รองรับการดำเนินการที่ไม่สลับที่ | เหมือนด้านบน แต่เธรดแต่ละตัวต้องเขียนลงหน่วยความจำร่วมเพื่อรับประกันลำดับการดำเนินการ |
| BLOCK_REDUCE_WARP_REDUCTIONS | กลยุทธ์รีดักชันระดับวาร์ปที่มีความหน่วงต่ำ | ① รีดักชันรีจิสเตอร์ → ② รีดักชัน Kogge-Stone ภายในวาร์ป → ③ การแพร่กระจายผลลัพธ์แบบอนุกรมระหว่างวาร์ป |
| BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC | รูปแบบที่ไม่แน่นอนซึ่งใช้การดำเนินการอะตอมมิกสำหรับการรวมระหว่างวาร์ป | ① รีดักชันรีจิสเตอร์ → ② รีดักชัน Kogge-Stone ภายในวาร์ป → ③ การสะสมอะตอมมิก |

การนำรูปแบบเหล่านี้ไปใช้อาศัยไฟล์เฉพาะสามไฟล์ ระบบจะเลือกเส้นทางที่

⚠️ หมายเหตุ: เนื้อหาได้รับการแปลโดย AI และตรวจสอบโดยมนุษย์ หากมีข้อผิดพลาดโปรดแจ้ง

☕ สนับสนุนค่ากาแฟทีมงาน

หากคุณชอบบทความนี้ สามารถสนับสนุนเราได้ผ่าน PromptPay

SCAN TO PAY WITH ANY BANK