SPMD เป็นหัวใจสำคัญของ OpenCL: งานย่อย (work-item) หลายพันรายการดำเนินการแบบขนานภายในโค้ดเคอร์เนลเดียวกัน และเมื่อพบ barrier() ก็จะซิงโครไนซ์กันอย่างเป็นระเบียบ

โมเดลการดำเนินการนี้ได้รับการสนับสนุนโดยตรงจากฮาร์ดแวร์บน GPU อย่างไรก็ตาม เมื่อแพลตฟอร์มเป้าหมายเปลี่ยนเป็น CPU ที่สามารถดำเนินการแบบเธรดเดียวตามลำดับ เช่น x86_64, ARM64 หรือ RISC-V ปัญหาก็เกิดขึ้น——จะทำให้ CPU คอร์เดียวจำลองงานย่อย 256 รายการพร้อมกัน และทำให้เกิดเอฟเฟกต์ “ทุกคนมาพร้อมกัน” ที่จุดกั้นได้อย่างไร?

นี่คือความท้าทายหลักที่ pocl/pocl[1] ต้องเอาชนะ

ในฐานะที่เป็นการใช้งาน OpenCL 3.0 แบบโอเพนซอร์สเพียงรายเดียวที่ผ่านการรับรองความสอดคล้องอย่างเป็นทางการจาก Khronos ทั้งบน CPU[2] และ Level Zero GPU[3] PoCL เสนอวิธีแก้ปัญหาดังนี้: ในระดับ LLVM IR ผ่าน Compiler Pass จะ “พับ” ความขนานของงานย่อยให้เป็น “ลูป” และเปลี่ยนสิ่งกีดขวางให้เป็นขอบเขตของลูป

บทความนี้จะเจาะลึกโค้ดกว่า 600,000 บรรทัด โดยมุ่งเน้นไปที่ lib/llvmopencl/ ซึ่งเป็น “เครื่องยนต์หลัก” ของ PoCL เพื่อเปิดเผยว่ามันเปลี่ยนเคอร์เนลที่เขียนขึ้นด้วยแนวคิดของ GPU ให้กลายเป็นรหัสเครื่องที่ CPU ก็สามารถดำเนินการได้อย่างมีประสิทธิภาพได้อย่างไร

เริ่มต้นใช้งานอย่างรวดเร็ว

ต่อไปนี้เป็นขั้นตอนที่ง่ายที่สุดในการสร้างและรันไดรเวอร์ CPU ของ PoCL บนระบบ Ubuntu สำหรับคำอธิบายการพึ่งพาโดยละเอียด การคอมไพล์ข้ามแพลตฟอร์ม และคำแนะนำในการสร้างแบ็กเอนด์ CUDA / Level Zero / Vulkan โปรดดูคู่มือการติดตั้งอย่างเป็นทางการ[4] และ README.md[5] ของที่เก็บ

# 1. ติดตั้ง dependencies (LLVM/Clang ต้องเป็นรุ่นล่าสุด)
sudo apt install -y cmake ninja-build python3 
libclang-dev libclang-cpp-dev libllvm-dev 
zlib1g-dev libtinfo-dev libhwloc-dev ocl-icd-libopencl1 ocl-icd-opencl-dev

# 2. ดึงโค้ดและ build (เปิดใช้งาน CPU driver + ICD โดยค่าเริ่มต้น)
git clone https://github.com/pocl/pocl.git && cd pocl
cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release
cmake --build build -j

# 3. สามารถใช้งานได้ทันทีโดยไม่ต้องติดตั้ง: ลงทะเบียน ICD จากไดเรกทอรี build ให้กับ OpenCL Loader
export OCL_ICD_VENDORS=$PWD/build/ocl-vendors
clinfo | head    # ควรเห็นแพลตฟอร์ม "Portable Computing Language"

หากเพียงต้องการทดลองใช้ แพ็คเกจไบนารีของดิสทริบิวชันก็ใช้งานได้เช่นกัน (เช่น apt install pocl-opencl-icd, brew install pocl, mamba install pocl-cuda)

สำหรับรายละเอียดเกี่ยวกับการทดสอบความสอดคล้อง CTS, การรองรับ SPIR-V ฯลฯ โปรดดูตัวเลือก เช่น ENABLE_CONFORMANCE, ENABLE_SPIRV ในเอกสารการสร้าง[6]

หนึ่ง ภาพรวมสถาปัตยกรรมและปรัชญาการออกแบบ

1.1 โครงสร้างสามชั้น: Host Runtime / Kernel Compiler / Device Driver

เมื่อเปิดไดเรกทอรีรากของที่เก็บ โครงสร้างทางกายภาพของ PoCL นั้นค่อนข้างเรียบง่าย ประกอบด้วยสามโมดูลหลัก:

PoCL ใช้ LLVM Pass เพื่อ “คอมไพล์” แนวคิด GPU ให้เป็นโค้ด CPU ได้อย่างไร: เวทมนตร์ SPMD ของ OpenCL 3.0 บน CPU

ส่วนที่ 2/6

รายละเอียดส่วนประกอบสถาปัตยกรรมหลัก

ฐานโค้ดของ PoCL จัดระเบียบรอบ ๆ สามไดเรกทอรีหลัก ซึ่งแต่ละแห่งมีหน้าที่แตกต่างกัน:

• lib/CL/: โฮสต์รันไทม์ (Host Runtime). ไฟล์ clXxx.c แต่ละไฟล์สอดคล้องกับจุดเข้าใช้งาน API ของ OpenCL เช่น clEnqueueNDRangeKernel.c[7] และ clCreateBuffer.c[8] งานหลักของไฟล์เหล่านี้รวมถึง: การตรวจสอบพารามิเตอร์, การจัดการการนับอ้างอิง, การบรรจุคำสั่ง และสุดท้ายมอบหมายงานให้กับเลเยอร์ไดรเวอร์อุปกรณ์ โครงสร้างข้อมูลที่ใช้ร่วมกันจะเน้นอยู่ในไฟล์ส่วนหัว pocl_cl.h[9] ที่มีขนาดใหญ่

• lib/llvmopencl/: คอมไพเลอร์เคอร์เนล (Kernel Compiler). นี่คือจิตวิญญาณที่ทำให้ PoCL แตกต่างจากโปรเจกต์ “LLVM tuning” ทั่วไป ไดเรกทอรีนี้ประกอบด้วย LLVM Function/Module Pass แบบกำหนดเองหลายสิบตัว ซึ่งมีหน้าที่แปลง IR เคอร์เนล OpenCL ที่สร้างโดย Clang ให้เป็นรูปแบบที่ “เวิร์กกรุ๊ปหนึ่งคือฟังก์ชันปกติหนึ่งฟังก์ชัน” ทำให้เคอร์เนลสามารถทำงานบน CPU ได้ด้วยลูปแบบอนุกรมควบคู่กับคำสั่ง SIMD

• lib/CL/devices/: ไดรเวอร์อุปกรณ์ (Device Drivers). แบ็กเอนด์ CPU, CUDA, Level Zero, Vulkan, Remote, OpenASIP ฯลฯ อยู่ในไดเรกทอรีนี้ โดยเชื่อมต่อกับระบบรันไทม์ผ่านตารางฟังก์ชัน pocl_devices_ops ที่เป็นหนึ่งเดียว

กระบวนการทั้งหมดสามารถสรุปเป็นห่วงโซ่การประมวลผลที่ “แคบทั้งสองด้าน กว้างตรงกลาง”:

clBuildProgram ─► Clang(ฟร้อนท์เอนด์) ─► LLVM IR ─► สายการผลิต Pass ของ PoCL ─► ISA อุปกรณ์ / SPIR-V ─► clEnqueueNDRangeKernel
(lib/llvmopencl) (lib/CL/devices/*)

1.2 เหตุใดจึงต้องมี Kernel Compiler เฉพาะ

เมื่อโปรแกรมเมอร์ OpenCL เขียน kernel void k(...) { ... barrier(CLK_LOCAL_MEM_FENCE); ... } สมมติฐานคือ: มีงานย่อย N = local_size_x × y × z รายการทำงานพร้อมกัน และซิงโครไนซ์กันที่ barrier()

• GPU รองรับความหมายนี้โดยธรรมชาติผ่านกลไกการจัดตาราง warp/wave;
• แต่ในสภาพแวดล้อม CPU หากเริ่มเธรดหนึ่งเธรดต่องานย่อยอย่างง่าย ๆ ค่า N มักจะสูงถึงหลายพัน และโอเวอร์เฮดของการจัดตารางเธรดจะบดบังการคำนวณจริง

ข้อมูลเชิงลึกที่สำคัญของ PoCL คือ: สิ่งกีดขวางจะแบ่งการควบคุมการไหลของเคอร์เนล——ส่วนของโค้ดระหว่างสิ่งกีดขวางสองอันนั้น “ดำเนินการแบบขนานโดยไม่รบกวนกัน” สำหรับงานย่อยทั้งหมด ซึ่งสามารถพับเป็นลูปอนุกรม/ที่สามารถทำเวกเตอร์ได้พอดี for (lid=0; lid<local_size; ++lid). ข้อมูลเชิงลึกนี้ถูกเสนอครั้งแรกโดย Pekka Jääskeläinen และคณะในเอกสารปี 2015 และ PoCL คือการใช้งานอ้างอิง

ต่อไปนี้ เราจะติดตาม IR เข้าสู่สายการผลิต Pass และสำรวจลงไปเรื่อย ๆ

สอง Kernel Compiler: ทำให้สิ่งกีดขวางกลายเป็นขอบเขตลูป

2.1 นามธรรมหลัก: Barrier และ ParallelRegion

PoCL จะทำให้การเรียก barrier() ของ OpenCL เป็นมาตรฐานก่อนเป็น CallInst LLVM พิเศษ——pocl.barrier

มันถูกแสดงโดยคลาสหุ้มบาง Barrier[10] ที่สืบทอดจาก llvm::CallInst:

// ที่มา: lib/llvmopencl/Barrier.h
class Barrier : public llvm::CallInst {
public:
// แทรกสิ่งกีดขวางที่จุดเริ่มต้น/สิ้นสุดของ BB; ใช้ซ้ำหากมีอยู่แล้ว
static Barrier *createAtStart(llvm::BasicBlock *BB) {
return create(BB->getFirstInsertionPt());
}
static Barrier *createAtEnd(llvm::BasicBlock *BB) {
return create(BB->getTerminator()->getIterator());
}
// ตรวจสอบว่าลูปมีสิ่งกีดขวางหรือไม่——การสอบถามที่สำคัญมาก
static bool isLoopWithBarrier(llvm::Loop &L) {
for (auto *BB : L.blocks())
for (auto &I : *BB)
if (llvm::isa<Barrier>(&I)) return true;
return false;
}
// ใช้ LLVM RTTI เพื่อให้ isa<Barrier>() ทำงาน
static bool classof(const llvm::CallInst *C) {
return C->getCalledFunction() &&
C->getCalledFunction()->getName() == BARRIER_FUNCTION_NAME;
}
};

เมื่อมี Barrier แล้ว PoCL สามารถดำเนินการทางทฤษฎีกราฟ: ตัดกราฟควบคุมการไหล (CFG) ตามความสัมพันธ์ “เส้นทางต้องผ่านสิ่งกีดขวางหนึ่งอัน” ออกเป็น ParallelRegion หลายอัน——นี่คือแนวคิดหลักที่สุดของคอมไพเลอร์ PoCL ซึ่งกำหนดไว้ใน ParallelRegion.h[11]

ParallelRegion คือกราฟย่อยที่เชื่อมต่อกันโดยไม่มีสิ่งกีดขวางผ่านในงานย่อยทั้งหมด โดยมีบล็อกพื้นฐานสิ่งกีดขวางอยู่ก่อนทางเข้าและหลังทางออก เมื่อแบ่งเสร็จแล้ว สำหรับ ParallelRegion เดียวกัน การดำเนินการของงานย่อยทั้งหมดจะเป็นอิสระต่อกัน——นี่คือพื้นฐานความถูกต้องสำหรับการนำไปใส่ในลูปงานย่อยในภายหลัง

2.2 การทำให้สิ่งกีดขวางเป็นมาตรฐาน: CanonicalizeBarriers และ ImplicitLoopBarriers

เพื่อให้การแบ่ง ParallelRegion เรียบร้อย PoCL จะดำเนินการ “จัดรูปแบบ” กับ IR ก่อน Pass ที่สำคัญที่สุดรวมถึง:

• CanonicalizeBarriers[12]: ตรวจสอบให้แน่ใจว่าสิ่งกีดขวางแต่ละอันครอบครองบล็อกพื้นฐานเดียว และบล็อกพื้นฐานนั้นมีเพียงคำสั่งสิ่งกีดขวางเท่านั้น ทำให้สิ่งกีดขวางกลายเป็น “ขอบเขตกราฟ” บริสุทธิ์
• ImplicitLoopBarriers[13]: หากลูปใดมีสิ่งกีดขวางอยู่ภายใน ตามความหมายของ OpenCL งานย่อยทั้งหมดต้องซิงโครไนซ์ในทุกการวนซ้ำของลูป Pass นี้จะแทรกสิ่งกีดขวางโดยนัยก่อนและหลังลูป เพื่อป้องกันไม่ให้การแบ่ง ParallelRegion ในภายหลังเข้าใจผิดว่าลูปเป็นพื้นที่ที่สามารถขนานกันอย่างอิสระ
• ImplicitConditionalBarriers[14]: จัดการกับสิ่งกีดขวางที่ “เส้นทางทั้งสองด้านของเงื่อนไขไม่สอดคล้องกัน” โดยยกระดับไปยังจุดก่อนการแตกกิ่งหรือจุดบรรจบ เพื่อหลีกเลี่ยงปัญหาการหยุดชะงักที่เกิดจากความแตกต่าง
• LoopBarriers[15]: ตรวจสอบให้แน่ใจว่า header และ latch ของลูปที่มีสิ่งกีดขวางต่างมีบล็อกพื้นฐานสิ่งกีดขวางของตัวเอง

Pass ก่อนการประมวลผลที่ดูเหมือน “จุกจิก” เหล่านี้ร่วมกันรักษาค่าคงที่: สิ่งกีดขวางที่แตกต่างกันทั้งหมดต้องมีอยู่เท่าเทียมกันบนเส้นทางการดำเนินการของงานย่อยทั้งหมด——มิฉะนั้น ความหมาย SPMD จะไม่สามารถพับเป็นลูปได้

2.3 กลยุทธ์การเลือก: WorkitemHandlerChooser

หลังจากแบ่ง ParallelRegion เสร็จ คำถามต่อไปคือ: จะดำเนินการ “ทำ ParallelRegion ซ้ำกับงานย่อยทั้งหมด” ได้อย่างไร?

PoCL กำหนดเส้นทางเทคนิคสองเส้นทางในไฟล์ WorkitemHandlerChooser.h[16]:

// ที่มา: lib/llvmopencl/WorkitemHandlerChooser.h  
enum class WorkitemHandlerType { LOOPS, CBS, INVALID };  

class WorkitemHandlerChooser  
: public llvm::AnalysisInfoMixin<WorkitemHandlerChooser> {  
public:  
using Result = WorkitemHandlerResult;     // แค่เก็บ enum  
Result run(llvm::Function &F, llvm::FunctionAnalysisManager &AM);  
};  

WorkitemHandlerType ChooseWorkitemHandler(llvm::Function &F);  

• LOOPS: ดำเนินการผ่าน Pass WorkitemLoops[17] หลักการคือการซ้อนลูป for 1 ถึง 3 มิติ (สอดคล้องกับ local_id ของมิติ X/Y/Z ตามลำดับ) ภายนอก ParallelRegion แต่ละอัน นี่คือกลยุทธ์เริ่มต้นและ成熟ที่สุดของระบบ
• CBS (Continuation-Based Synchronization): ดำเนินการโดย SubCFGFormation[18] แนวทางนี้ใช้กลไกการคัดลอก sub-CFG และการสลับสถานะ ซึ่งทำงานได้ดีกว่าเมื่อจัดการกับเคอร์เนลที่มีการควบคุมการไหลที่ไม่สม่ำเสมอ และยังเอื้อต่อความสามารถในการเพิ่มประสิทธิภาพของ LoopVectorize ของ LLVM มากกว่า

Chooser ในฐานะส่วนประกอบ LLVM Analysis จะไม่แก้ไข IR หน้าที่ของมันคือการแนบผลการตัดสินใจเข้ากับตัวจัดการการวิเคราะห์ เพื่อให้ Pass อื่น ๆ สามารถนำไปใช้ได้ตามต้องการ

2.4 ตัวเอก登场: WorkitemLoops ห่อหุ้มลูปอย่างไร

คลาสการใช้งาน WorkitemLoopsImpl[19] ของ WorkitemLoops ภายในประกอบด้วย DominatorTree, LoopInfo, PostDominatorTree ของ LLVM และ VariableUniformityAnalysis ที่พัฒนาเองของ PoCL

การดำเนินการหลักคือ: สำหรับ ParallelRegion แต่ละอัน ให้เรียก createLoopAround เพื่อห่อหุ้มลูปหนึ่งชั้นตามลำดับในมิติ X, Y, Z

// ที่มา: lib/llvmopencl/WorkitemLoops.cc  
class WorkitemLoopsImpl :public pocl::WorkitemHandler {  
// สถานะสำคัญ  
llvm::DominatorTree &DT;  
llvm::LoopInfo &LI;  
llvm::PostDominatorTree &PDT;  
VariableUniformityAnalysisResult &VUA;         // ทำเครื่องหมายว่าตัวแปรใดมีค่าเท่ากันสำหรับ WI ทั้งหมด  
ParallelRegion::ParallelRegionVector OriginalParallelRegions; // ผลการแบ่ง  
StrInstructionMap ContextArrays;               // "อาร์เรย์บริบท" สำหรับตัวแปรข้ามภูมิภาค  
std::array<llvm::GlobalVariable *, 3> GlobalIdIterators;  
llvm::Value *LocalIdXFirstVar;                 // จุดเริ่มต้นการวนซ้ำครั้งแรก (ใช้สำหรับ 0,0,0 ที่ถูก peel แล้ว)  

// สร้างลูป for ระหว่าง [EntryBB, ExitBB] รอบ Region ในมิติ Dim  
std::pair<llvm::BasicBlock *, llvm::BasicBlock *>  
createLoopAround(ParallelRegion &Region,  
llvm::BasicBlock *EntryBB,  
llvm::BasicBlock *ExitBB,  
bool PeeledFirst, int Dim,  
bool AddIncBlock = true,  
llvm::Value *DynamicLocalSize = nullptr);  
};  

เทมเพลต IR ที่สร้างขึ้นภายในโดย createLoopAround (เนื้อหาต่อไปนี้แปลจากคำอธิบายประกอบภาษาอังกฤษในไฟล์ต้นฉบับ) มีโครงสร้างประมาณดังนี้:

; ที่มา: เทมเพลตคำอธิบายประกอบ createLoopAround ใน lib/llvmopencl/WorkitemLoops.cc  
for.init:  
store i32 0, i32* %_local_id_x        ; เริ่มต้นปกติเมื่อ PeeledFirst=false  
br label %for.cond  
for.cond:  
%lid = load i32, i32* %_local_id_x  
%cmp = icmp ult i32 %lid, %local_size_x  
br i1 %cmp, label %region.entry, label %for.end  
region.entry:  
... ; โค้ดเดิมของ ParallelRegion  
br label %for.inc  
for.inc:  
%next = add i32 %lid, 1  
store i32 %next, i32* %_local_id_x  
br label %for.cond  
for.end:                                ; คือ ExitBB ชั้นนอก / BB ที่มีสิ่งกีดขวาง  

หลังจากซ้อนลูปที่คล้ายกันอีกหนึ่งชั้นในมิติ Y และ Z การดำเนินการของงานย่อย N = X·Y·Z ครั้งภายในเวิร์กกรุ๊ปทั้งหมดจะถูกเชื่อมต่อเป็นลูปตามลำดับ เมื่อ LLVM รู้จัก “ลูปตัวแปรเหนี่ยวนำมาตรฐาน” นี้แล้ว มันจะสามารถเปิดใช้งานการเพิ่มประสิทธิภาพทั่วไป เช่น LoopVectorize, SLP vectorization, Loop Unroll ได้——นี่คือเคล็ดลับหลักที่ทำให้ PoCL ทำให้เคอร์เนล OpenCL ทำงานบน CPU ได้อย่างมีประสิทธิภาพ

2.5 “อาร์เรย์บริบท” สำหรับตัวแปรข้ามภูมิภาค

แค่ห่อหุ้มลูปอย่างเดียวไม่พอ พิจารณาโค้ดที่ลดรูปแล้วนี้:

int x = compute();  
barrier(CLK_LOCAL_MEM_FENCE);  
use(x);  

ตัวแปร x ถูกกำหนดก่อนสิ่งกีดขวางและถูกใช้หลังจากสิ่งกีดขวาง——ซึ่งหมายความว่าสิ่งกีดขวางแบ่งโค้ดออกเป็นสอง ParallelRegion และ ค่า x ของแต่ละงานย่อยจะต้องถูกเก็บรักษาไว้อย่างสมบูรณ์ จนถึง Region ที่สองจึงจะสามารถใช้งานได้อย่างถูกต้อง

วิธีแก้ปัญหาของ PoCL คือ: ยกระดับ “ตัวแปรที่มีชีวิตข้ามภูมิภาค” นี้เป็นอาร์เรย์ alloca ที่มีความยาวเท่ากับ local_size_x*y*z (นั่นคือ ContextArrays ในโค้ด) การดำเนินการเฉพาะคือ: ภายในลูปของ Region แรก ให้เก็บค่าตามดัชนี local_id ลงในอาร์เรย์; ภายในลูปของ Region ที่สอง ให้นำค่าออกมาตามดัชนี local_id เดียวกัน

WorkitemLoopsImpl::fixMultiRegionVariables คือฟังก์ชันที่รับผิดชอบตรรกะนี้ มันใช้ประโยชน์จากแผนผังการครอบงำและแผนผังการครอบงำภายหลังที่สร้างไว้ล่วงหน้าในส่วนก่อนหน้า รวมถึงผลการวิเคราะห์ของ VariableUniformityAnalysis——ซึ่งสามารถระบุตัวแปรที่ “มีค่าเท่ากันสำหรับงานย่อยทั้งหมด” (เช่น ค่าคงที่ในลูปและพารามิเตอร์เคอร์เนล) เพื่อหลีกเลี่ยงการจัดเก็บบริบทที่ไม่จำเป็นสำหรับตัวแปรเหล่านี้ ช่วยประหยัดพื้นที่สแต็กและเวลาในการดำเนินการ

2.6 เครื่องมือช่วยเหลือ เช่น PHIsToAllocas, IsolateRegions

เพื่อให้แน่ใจว่าการจัดเก็บบริบทสามารถดำเนินการได้อย่างถูกต้อง PoCL ยังมี Pass ที่ “ดูไม่เด่น แต่ขาดไม่ได้” อีกหลายตัว:

• PHIsToAllocas[20]: ลดระดับโหนด PHI เป็นการดำเนินการอ่าน/เขียนหน่วยความจำ เพื่อให้ง่ายต่อการจัดการวงจรชีวิตของตัวแปรผ่านอาร์เรย์ alloca ภายนอกลูป
• IsolateRegions[21]: แทรกบล็อกพื้นฐานว่างที่ขอบเขตของ Region เพื่อป้องกันไม่ให้ขอบระหว่าง Region ต่าง ๆ รบกวนกัน
• AllocasToEntry[22]: ย้ายคำสั่ง alloca ทั้งหมดไปยังบล็อกทางเข้า เพื่อสร้างความสะดวกสำหรับการเพิ่มประสิทธิภาพ mem2reg ในภายหลัง
• BarrierTailReplication[23]: จัดการกับโค้ด “ส่วนท้าย” หลังจากสิ่งกีดขวาง โดยคัดลอกในหลายบล็อกก่อนหน้า เพื่อสร้างแบบจำลองความหมายหลายทางเข้าของสิ่งกีดขวางอย่างถูกต้อง

Pass เหล่านี้ทำงานร่วมกันเพื่อปรับรูปร่าง LLVM IR ให้เป็นรูปแบบที่ “เป็นมิตรกับลูปงานย่อย”

2.7 จุดสิ้นสุด: การห่อหุ้มฟังก์ชัน Workgroup

หลังจากการแปลงชุดข้างต้น แต่ละเคอร์เนล OpenCL K จะถูกคอมไพล์เป็นฟังก์ชัน C ที่เทียบเท่า _pocl_kernel_K_workgroup ซึ่งลายเซ็นประมาณดังนี้:

// ที่มา: lib/llvmopencl/Workgroup.h (ลายเซ็นแนวคิด)
void _pocl_kernel_K_workgroup(uint8_t **args,
void *pocl_context,    // group_id, local_size, ...
uint32_t group_x,
uint32_t group_y,
uint32_t group_z);

Workgroup.cc[24] มีหน้าที่สร้างโค้ดห่อหุ้มชั้นนี้: มัน “แกะ” พารามิเตอร์ของเคอร์เนล OpenCL จากอาร์เรย์พอยน์เตอร์แบบรวม จากนั้นเรียกเนื้อหาเคอร์เนลที่มีลูปงานย่อยอยู่แล้ว ไดรเวอร์อุปกรณ์เพียงแค่ต้องต่อลูป for (group_x...) for (group_y...) for (group_z...) workgroup(...) ทางฝั่งโฮสต์ ในขณะที่ความขนานระดับงานย่อยถูกจัดการอย่างสมบูรณ์ในขั้นตอนการคอมไพล์

สาม Host Runtime: จาก API สู่คำสั่งอุปกรณ์

3.1 ICD และ Platform Dispatch

โดยทั่วไปแล้วแอปพลิเคชัน OpenCL จะเรียกใช้การใช้งานเฉพาะของผู้ผลิตผ่าน ICD Loader

PoCL ใช้จุดเข้า ICD ใน clIcdGetFunctionAddressForPlatformKHR.c[25] และ clIcdGetPlatformIDsKHR.c[26]; pocl_intfn.h[27] จะรวบรวมฟังก์ชันเข้า clXxx ทั้งหมด และแปลงเป็นพอยน์เตอร์ฟังก์ชันภายในที่สามารถจัดทำดัชนีผ่านตาราง dispatch

3.2 กลไกคิวคำสั่งและอีเวนต์

OpenCL ใช้ความหมายแบบอะซิงโครนัส: ฟังก์ชันชุด clEnqueueXxx วางคำสั่งในคิว และใช้วัตถุ cl_event เพื่อเชื่อมโยงการพึ่งพา

PoCL ใช้กลไกหลักต่อไปนี้ใน pocl_util.c[28] (73 KB, ถือเป็น “มีดพับสวิส” ของเลเยอร์ Runtime) และ pocl_mem_management.c[29]:

• pocl_create_event, pocl_create_command: ใช้สำหรับสร้างโหนดคำสั่งและวัตถุอีเวนต์
• เครื่องสถานะอีเวนต์: เป็นไปตามการเปลี่ยนสถานะ SUBMITTED → READY → RUNNING → COMPLETE และต้องจัดการการย้ายข้อมูลเมื่อข้ามอุปกรณ์
• ตรรกะการจัดตาราง: วางคำสั่งในฮุค push_command ที่ไดรเวอร์อุปกรณ์จัดเตรียมไว้

clEnqueueNDRangeKernel.c มีหน้าที่หลักในการตรวจสอบพารามิเตอร์ จากนั้นส่งต่องานให้กับ pocl_ndrange_kernel.c ส่วนหลังจะดำเนินการแบ่ง work-group, แก้ไข work_dim, ลงทะเบียนพอยน์เตอร์ SVM/USM ฯลฯ และสุดท้ายสร้างคำสั่งประเภท CL_COMMAND_NDRANGE_KERNEL และส่งไปยังคิวอุปกรณ์

3.3 การสร้างโปรแกรม: จากซอร์สโค้ด/SPIR-V สู่ LLVM bitcode

clBuildProgram ดูเหมือนเป็นฟังก์ชันห่อหุ้มที่เรียบง่าย แต่งานหนักจริง ๆ ดำเนินการโดย pocl_build.c[30] และ pocl_llvm_build.cc[31] (41 KB)

pocl_llvm_build.cc เรียกใช้ C++ API ของ Clang โดยตรง เพื่อทำการคอมไพล์ฟร้อนท์เอนด์จากซอร์ส OpenCL C เป็น LLVM IR pocl_llvm_spirv.cc[32] จัดการกับอินพุตและเอาต์พุตในรูปแบบ SPIR-V bitcode ที่สร้างขึ้นจะถูกแคชผ่าน pocl_cache.c[33] โดยใช้แฮชตามเนื้อหา เพื่อหลีกเลี่ยงการสร้างซ้ำ——นี่คือสาเหตุที่ PoCL สามารถ “อุ่นเครื่อง” และเร่งความเร็วได้อย่างมากหลังจากแอปพลิเคชันขนาดใหญ่เริ่มทำงานครั้งแรก

สี่ Device Driver: การรองรับหลายแบ็กเอนด์ภายใต้อินเทอร์เฟซแบบรวม

4.1 ตารางฟังก์ชัน pocl_devices_ops

ไดรเวอร์อุปกรณ์แต่ละตัวเชื่อมต่อกับ Runtime โดยใช้ชุดฟังก์ชัน callback ที่มีโครงสร้างเหมือนกัน: รวมถึงการจัดสรรหน่วยความจำ, การจัดตารางคำสั่ง, การดำเนินการเคอร์เนล, การเข้าถึงอิมเมจ ฯลฯ

• ไดรเวอร์ CPU อยู่ในไดเรกทอรีย่อย เช่น lib/CL/devices/pthread/, lib/CL/devices/basic/
• แบ็กเอนด์ GPU อยู่ในไดเรกทอรี เช่น lib/CL/devices/level0/, lib/CL/devices/cuda/

การออกแบบ “ไดรเวอร์แบบบาง + Runtime แบบหนา” นี้ทำให้การทดสอบความสอดคล้อง OpenCL (CTS) เพียงแค่แก้ไขบั๊กครั้งเดียวในเลเยอร์ Runtime ไดรเวอร์ทั้งหมดก็จะได้รับประโยชน์โดยอัตโนมัติ

4.2 ไดรเวอร์ CPU: เรียกใช้ฟังก์ชัน Workgroup เป็นพอยน์เตอร์ฟังก์ชัน

สำหรับแต่ละคำสั่ง NDRange ไดรเวอร์ CPU จะดำเนินการลูป group สามชั้นทางฝั่งโฮสต์ โดยเรียก workgroup_fn(args, &ctx, gx, gy, gz) โดยตรงในแต่ละครั้ง

เนื่องจาก workgroup_fn มีลูปงานย่อยและโค้ดท้องถิ่นที่ทำเวกเตอร์โดย LLVM อยู่แล้ว ประสิทธิภาพโดยรวมเมื่อจัดการกับเคอร์เนลมาตรฐาน เช่น BLAS, convolution, reduce สามารถเข้าใกล้ระดับของการใช้งาน OpenMP ที่เขียนด้วยมือ

นอกจากนี้ ไดรเวอร์ CPU ยังใช้ hwloc[34] เพื่อระบุโทโพโลยี CPU โดยอัตโนมัติ จัดตาราง group ไปยังคอร์ที่เหมาะสม และใช้พูลเธรดน้ำหนักเบาที่ implement ใน pocl_threads.c[35] เพื่อลดความกระวนกระวายในการจัดตาราง

4.3 ไดรเวอร์ GPU: การสลับเป้าหมายในขั้นตอนการคอมไพล์

สำหรับแบ็กเอนด์ CUDA และ Level Zero สายการผลิต Pass ชุดเดียวกันใน lib/llvmopencl/ ก็ใช้ได้เช่นกัน เพียงแต่เปลี่ยนแพลตฟอร์มเป้าหมายของ LLVM ในขั้นตอนการสร้าง Workgroup.cc[36] และ SPIR-V/PTX——กลยุทธ์ “ฟร้อนท์เอนด์และมิดเดิลเอนด์ทั่วไป แบ็กเอนด์แยกตามอุปกรณ์” นี้คือสาเหตุหลักที่ PoCL สามารถครอบคลุมแพลตฟอร์มฮาร์ดแวร์ได้มากมายด้วยโค้ดชุดเดียว

ห้า รายละเอียดทางวิศวกรรมที่น่าจดจำ

5.1 การออกแบบ Pass ที่ขับเคลื่อนด้วยค่าคงที่

ย้อนดูการออกแบบ Pass ภายใต้ lib/llvmopencl/ เราจะเห็นแนวคิดแบบแบ่งชั้นที่เรียบง่ายแต่มีประสิทธิภาพ: ขั้นแรก รักษาค่าคงที่ทางความหมายผ่าน Pass ก่อนการประมวลผล (เช่น สิ่งกีดขวางเป็นเส้นทางบังคับสำหรับทุกงานย่อย ลูปที่มีสิ่งกีดขวางต้องซิงโครไนซ์) จากนั้นให้ Pass ที่แก้ไขโค้ดจริง (เช่น WorkitemLoops / SubCFGFormation) ทำงานบน IR ที่สะอาด ข้อดีของการออกแบบนี้คือ ความซับซ้อนของแต่ละ Pass สามารถควบคุมได้ และสามารถทดสอบแยกกัน——ซึ่งสำคัญมากสำหรับโปรเจกต์ที่มีโค้ดกว่า 600,000 บรรทัดและต้องรองรับ LLVM หลายรุ่นใหญ่ตั้งแต่ 18 ถึง 22

5.2 การผูกพันอย่างแน่นหนากับระบบนิเวศต้นน้ำของ LLVM

PoCL เชื่อมต่อกับ LLVM internal API อย่างแน่นหนาผ่าน #include <llvm/...> ซึ่งเป็นดาบสองคม: ด้านหนึ่ง มันสามารถใช้ประโยชน์จากเทคนิคการเพิ่มประสิทธิภาพล่าสุดได้อย่างเต็มที่ (เช่น การใช้ MinLegalVecSize[37] เพื่อปรับเกณฑ์การทำเวกเตอร์ และการจัดการพฤติกรรมที่ไม่ได้กำหนดผ่าน SanitizeUBofDivRem[38]); แต่อีกด้านหนึ่ง มันต้องปรับตัวเข้ากับการเปลี่ยนแปลงรุ่นใหญ่ของ LLVM ทุกปี——คำอธิบายความเข้ากันได้ของ Clang Bug สำหรับ CUDA/LLVM 21 ใน README เป็นตัวอย่างทั่วไป

5.3 ต้นทุนเบื้องหลังความสอดคล้อง

ข้อมูลอัตราการผ่าน CTS ที่ให้ไว้ใน README แสดงให้เห็นว่า: ไดรเวอร์ CPU/x86_64 สามารถผ่านการทดสอบ OpenCL 3.0 CTS ได้ 100% และไดรเวอร์ CPU/RISCV ก็ผ่านมากกว่า 99% เบื้องหลังผลลัพธ์ที่โดดเด่นเหล่านี้คือการแก้ไขและปรับแต่งกรณีขอบจำนวนมาก เช่น: ตรรกะการเริ่มต้นของ

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

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

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

SCAN TO PAY WITH ANY BANK