FlashRT: เอ็นจิ้นการอนุมานที่ออกแบบมาเพื่อการควบคุมหุ่นยนต์แบบเรียลไทม์โดยเฉพาะ หน่วงเวลาเพียง 2.39ms ประสิทธิภาพเหนือกว่า TensorRT

4 days ago • การอนุมานโมเดลขนาดใหญ่ • 59 views

การอภิปรายทั่วไปเกี่ยวกับการอนุมานโมเดลขนาดใหญ่มักจะมุ่งเน้นไปที่ “ปริมาณงานที่สูงขึ้น” “ขนาดแบตช์ที่ใหญ่ขึ้น” และ “ระบบบริการที่ซับซ้อนมากขึ้น”

อย่างไรก็ตาม ในสถานการณ์ต่างๆ เช่น การควบคุมหุ่นยนต์ การโต้ตอบแบบเรียลไทม์ และการปรับใช้บนขอบ (edge deployment) เรากำลังเผชิญกับความท้าทายที่แตกต่างกันโดยสิ้นเชิง: ขนาดแบตช์ที่เล็กมาก ความหน่วงที่ไวต่อความรู้สึกอย่างยิ่ง รูปแบบอินพุตที่เปลี่ยนแปลงบ่อยครั้ง และการควบคุมแบบวงปิดที่ไม่สามารถทนต่อค่าใช้จ่ายในการจัดตารางงานเพิ่มเติมอีกหลายสิบมิลลิวินาทีได้

TensorRT เชี่ยวชาญในการคอมไพล์โมเดลเป็นเอนจินที่ถูกแช่แข็ง (frozen engine) ในขณะที่ vLLM และ SGLang มุ่งเน้นไปที่การให้บริการ LLM ที่มีพร้อมกันสูง แต่เมื่อภารกิจเปลี่ยนเป็น “ป้อนภาพจากกล้อง รับการเคลื่อนไหวของแขนกลทันที” ตรรกะพื้นฐานของระบบอนุมานก็เปลี่ยนไปแล้ว

FlashRT เป็นเอนจินการอนุมานแบบเรียลไทม์ประสิทธิภาพสูงที่ออกแบบมาสำหรับปริมาณงาน AI ที่มีความหน่วงต่ำและขนาดแบตช์เล็กโดยเฉพาะ การประยุกต์ใช้หลักแบบบูรณาการคือการจัดหาฟังก์ชันการควบคุม VLA (Vision-Language-Action) ระดับการผลิตสำหรับ Pi0, Pi0.5, GROOT N1.6 และ Pi0-FAST รวมถึงรองรับโมเดลภาษาขนาดใหญ่ เช่น Qwen 3.6-27B พารามิเตอร์
ที่เก็บโค้ด: https://github.com/LiangSu8899/FlashRT
บทความเต็มประมาณ 9,000 คำ คาดว่าใช้เวลาอ่าน 40 นาที เวอร์ชันพอดแคสต์ประมาณ 30 นาที

FlashRT คือแนวทางปฏิบัติทางวิศวกรรมที่เติบโตมาจากเส้นทางที่แตกต่างนี้: มันไม่ได้พยายามมอบหมายงานทั้งหมดให้กับคอมไพเลอร์เอนกประสงค์ แต่จะแยกกระบวนการอนุมานไปข้างหน้าของโมเดลออกเป็น CUDA kernel, โปรโตคอลพอยน์เตอร์, น้ำหนักเชิงปริมาณ และ CUDA Graph แบบคงที่ที่เสถียร และสร้างช่องทางความเร็วสูงระดับต่ำที่มุ่งเน้นการควบคุมแบบเรียลไทม์ภายใต้ส่วนหน้าของ PyTorch/JAX

ตารางด้านล่างเปรียบเทียบความหน่วงและปริมาณงานของการอนุมานของโมเดล AI สี่รุ่น ได้แก่ Pi0.5, Pi0, GROOT N1.6 และ Pi0-FAST เวอร์ชันปรับปรุงประสิทธิภาพ บน Jetson AGX Thor (ฝั่งขอบ) และ RTX 5090 (ฝั่งเดสก์ท็อป) บนแพลตฟอร์ม Jetson ความหน่วงของโมเดลทั่วไปอยู่ที่ประมาณ 41-46ms ปริมาณงานอยู่ระหว่าง 22-24Hz ในขณะที่ Pi0-FAST หลังจากปรับปรุงกระบวนการแล้ว ความหน่วงลดลงเหลือ 8.1ms/token และปริมาณงานสูงถึง 123 tok/s บนแพลตฟอร์ม RTX 5090 ประสิทธิภาพของแต่ละโมเดลเพิ่มขึ้นอย่างมาก: Pi0 มีความหน่วงต่ำถึง 17.58-24.48ms, GROOT มีปริมาณงานสูงสุดถึง 80Hz, และ Pi0-FAST ลดลงเหลือ 2.39ms/token โดยมีปริมาณงานสูงถึง 418 tok/s สิ่งนี้แสดงให้เห็นถึงข้อได้เปรียบด้านพลังการคำนวณของ GPU เดสก์ท็อประดับสูง และยังยืนยันถึงคุณค่าที่สำคัญของการเพิ่มประสิทธิภาพการอนุมานสำหรับการปรับใช้ฝั่งปลายทาง (edge deployment) ในขณะเดียวกัน ความแตกต่างด้านประสิทธิภาพในสถานการณ์การสร้างมุมมองหลายมุมและหลาย token ยังสะท้อนให้เห็นถึงความสัมพันธ์ที่แน่นแฟ้นระหว่างความซับซ้อนของงานและความต้องการพลังการคำนวณ ซึ่งเป็นข้อมูลอ้างอิงที่ชัดเจนสำหรับการปรับใช้โมเดลในสถานการณ์ต่างๆ เช่น หุ่นยนต์และระบบฝังตัว

การเปรียบเทียบอีกชุดหนึ่งครอบคลุมโซลูชันการอนุมาน AI หลายแบบ เช่น Original openpi, PyTorch naive, torch.compile, Triton-based VLA, NVIDIA VLA-Perf, Isaac GR00T และ FlashRT บนฮาร์ดแวร์ Jetson Thor, RTX 4090/5090 สำหรับโมเดล Pi0, Pi0.5, GROOT N1.6 Original openpi ที่ไม่ได้ปรับแต่งมีความหน่วงสูงถึง 714ms บน Jetson Thor ในขณะที่ PyTorch naive อยู่ที่ประมาณ 200ms บน RTX 4090 หลังจากปรับแต่งด้วย torch.compile, Triton, TensorRT แล้ว ความหน่วงลดลงอย่างมีนัยสำคัญ โซลูชัน FlashRT ในครั้งนี้มีประสิทธิภาพดีที่สุด: บน RTX 5090 ความหน่วงของ Pi0, Pi0.5, GROOT N1.6 ต่ำถึง 21.16ms, 17.58ms และ 13.08ms ตามลำดับ บน Jetson Thor ก็ทำได้ต่ำถึง 46ms, 39.78-51.51ms และ 45ms เช่นกัน สิ่งนี้ไม่เพียงสะท้อนถึงคุณค่าที่สำคัญของการเพิ่มประสิทธิภาพการอนุมานสำหรับการปรับใช้บนขอบและเดสก์ท็อป แต่ยังเน้นย้ำถึงข้อได้เปรียบด้านประสิทธิภาพของ FlashRT บน GPU ฝั่งปลายทางและเดสก์ท็อป ซึ่งเป็นโซลูชันที่มีประสิทธิภาพสำหรับการปรับใช้แบบเรียลไทม์ในสถานการณ์ต่างๆ เช่น หุ่นยนต์

unsetunsetสารบัญunsetunset

หนึ่ง. เริ่มต้นใช้งานอย่างรวดเร็ว: รันให้ได้ก่อน แล้วค่อยเข้าใจว่าทำไมมันถึงเร็ว
สอง. FlashRT แก้ปัญหาอะไรได้บ้าง
- 2.1 มันไม่ใช่ LLM Serving ทั่วไป แต่เป็นสแต็กการอนุมานเพื่อการควบคุมแบบเรียลไทม์
- 2.2 ข้อมูลประสิทธิภาพหลักบ่งบอกถึงขอบเขตเป้าหมายของมัน
สาม. สถาปัตยกรรมโดยรวม: แบ่งการส่งต่อโมเดลออกเป็นสัญญาสี่ชั้น
- 3.1 ชั้น API สาธารณะ: ผู้ใช้เห็นแค่ load_model และ predict
- 3.2 ชั้น Frontend: อินพุต น้ำหนัก และการปรับเทียบที่เกี่ยวข้องกับเฟรมเวิร์ก
- 3.3 ชั้น Pipeline: การส่งต่อแบบพอยน์เตอร์ที่ไม่ขึ้นกับเฟรมเวิร์ก
- 3.4 ชั้น Hardware: attention และ GEMM จ่ายงานตามฮาร์ดแวร์
สี่. ทำไม CUDA Graph ถึงเป็นแกนหลักด้านประสิทธิภาพของ FlashRT
- 4.1 ในสถานการณ์ batch เล็ก การจัดตารางงานของ Python เองคือคอขวด
- 4.2 ข้อกำหนดในการจับภาพกราฟ: พอยน์เตอร์คงที่ ไม่มีการจัดสรรแบบไดนามิก รูปร่างคงที่
ห้า. การโหลดน้ำหนัก: เปลี่ยน state_dict ให้เป็นตารางข้อกำหนดแบบประกาศ
- 5.1 WEIGHT_SPEC คือคำอธิบายการคอมไพล์แบบน้ำหนักเบาของ FlashRT
- 5.2 ความแตกต่างระหว่าง Pi0 และ Pi0.5 ถูกเข้ารหัสไว้ใน spec ไม่ใช่กระจายอยู่ในตรรกะการอนุมาน
หก. AttentionBackend: ทำให้โมเดลไม่ต้องสนใจว่าเป็น FlashAttention หรือ Thor FMHA
- 6.1 pipeline เดียวกัน เบื้องหลังอาจเป็นการ implement attention ที่แตกต่างกัน
- 6.2 attention site ของ GROOT แสดงให้เห็นถึงวิธีการแยกส่วนโมเดลที่ซับซ้อน
เจ็ด. FP8 และ NVFP4: การหาปริมาณไม่ใช่การบีบอัด แต่เป็นการควบคุมการแพร่กระจายของข้อผิดพลาด
- 7.1 เส้นทาง FP8: การหาปริมาณน้ำหนักและการปรับเทียบแอคติเวชันเข้าสู่การจับภาพกราฟ
- 7.2 เส้นทาง NVFP4: ความพยายามที่รุนแรงยิ่งขึ้นสำหรับ Pi0.5 encoder FFN
แปด. เส้นทาง Qwen3.6: FlashRT กำลังขยายประสบการณ์ VLA ไปยัง LLM
- 8.1 แผนการย้ายข้อมูลแบบเป็นระยะของ Qwen3.6
- 8.2 เค้าโครงพอยน์เตอร์ของ KV cache ก็เพื่อความเสถียรของกราฟเช่นกัน
เก้า. การตอบสนองการฝึกอบรม: นำ inference kernel มาใช้โดยตรงในการฝึก forward
- 9.1 หลักการออกแบบ: inference kernel = training kernel
- 9.2 ข้อมูล: ความเร็วในการฝึกและ VRAM ลดลงทั้งคู่
- 9.3 ความหมายของวงปิด: แก้ปัญหาจุดเจ็บปวดของ “การปรับแต่งบ่อยครั้ง” ใน embodied intelligence
สิบ. ความแตกต่างพื้นฐานกับ TensorRT, vLLM, SGLang
- 10.1 FlashRT ปฏิเสธเส้นทางหนัก “ส่งออกก่อน คอมไพล์ก่อน แล้วค่อยปรับใช้”
- 10.2 มันไม่ใช่ตัวจัดตารางงานพร้อมกันสูงแบบ vLLM
สิบเอ็ด. การแลกเปลี่ยนทางวิศวกรรม: ทำไม FlashRT ถึงควรค่าแก่การศึกษา
- 11.1 ข้อดี: ทำให้ความหน่วงแบบ end-to-end เป็นเป้าหมายแรก
- 11.2 ต้นทุน: การปรับโมเดลต้องใช้ “การแปล” ทางวิศวกรรม
- 11.3 แรงบันดาลใจสูงสุด: AI แบบเรียลไทม์ต้องการปรัชญาระบบการอนุมานอีกแบบหนึ่ง
บทสรุป: ส่งคืนการปรับใช้โมเดลให้กับวงควบคุม

unsetunsetหนึ่ง. เริ่มต้นใช้งานอย่างรวดเร็ว: รันให้ได้ก่อน แล้วค่อยเข้าใจว่าทำไมมันถึงเร็วunsetunset

ตามเอกสาร README ของ FlashRT วิธีการเรียกใช้ที่ง่ายที่สุดต้องใช้ API เพียงสามบรรทัด: โหลดโมเดล ป้อนภาพและคำสั่งภาษา รับเอาต์พุตการกระทำ แต่โปรดทราบว่าโปรเจกต์นี้ขึ้นอยู่กับไฟล์ kernel .so ที่คอมไพล์ไว้แล้ว หากคุณไม่ได้ใช้สภาพแวดล้อมการ build ที่กำหนดค่าไว้ล่วงหน้า คุณต้องดำเนินการขั้นตอนการ build ด้วย CMake ก่อน

ต่อไปนี้คือตัวอย่างการเรียกใช้ Python ที่ง่ายที่สุด:

import flash_rt

model = flash_rt.load_model(
checkpoint="/path/to/pi05_checkpoint",
config="pi05",          # หรือ "pi0", "groot", "pi0fast"
framework="torch",      # หรือ "jax"
)

actions = model.predict(
images=[base_img, wrist_img],
prompt="pick up the red block",
)

หากคุณวางแผนที่จะใช้ Docker คำสั่ง build image ที่แนะนำใน README มีดังนี้:

docker pull ghcr.io/liangsu8899/flashrt:latest
docker run --rm --gpus all -it ghcr.io/liangsu8899/flashrt:latest

อย่างไรก็ตาม docker/README.md ในที่เก็บโค้ดก็ระบุเช่นกันว่าเนื่องจากชื่อแพ็คเกจถูก refactor จาก flash_vla เป็น flash_rt image สาธารณะอาจยังไม่ได้รับการอัปเดต ดังนั้น วิธีที่ปลอดภัยกว่าคือการ build ด้วยตัวเองในเครื่อง:

git clone https://github.com/LiangSu8899/FlashRT.git
cd FlashRT
docker build -t flashrt:dev -f docker/Dockerfile .
docker run --rm --gpus all -it flashrt:dev

ในสภาพแวดล้อม Linux ดั้งเดิม คำสั่ง build ที่จำเป็นสามารถสรุปได้ดังนี้:

python3.12 -m venv .venv
source .venv/bin/activate

pip install torch --index-url https://download.pytorch.org/whl/cu128
pip install pybind11 cmake "numpy>=1.24" safetensors
pip install "transformers<4.56" pandas pillow pyarrow

git clone https://github.com/LiangSu8899/FlashRT.git
cd FlashRT
git clone --depth 1 --branch v4.4.2 https://github.com/NVIDIA/cutlass.git third_party/cutlass

pip install -e ".[torch]"
cmake -B build -S .
cmake --build build -j$(nproc)

ตรวจสอบว่า build สำเร็จหรือไม่โดยรันคำสั่งต่อไปนี้:

python examples/quickstart.py 
    --checkpoint /path/to/pi05_checkpoint 
    --benchmark 20

สำหรับข้อมูลการพึ่งพาที่สมบูรณ์ยิ่งขึ้น ขอบเขตการครอบคลุมของสถาปัตยกรรม GPU และรายละเอียดการใช้งาน Docker โปรดดู README.md, docs/INSTALL.md และ docker/README.md ในโปรเจกต์

สอง. FlashRT แก้ปัญหาอะไรได้บ้าง

2.1 มันไม่ใช่เอนจินบริการโมเดลขนาดใหญ่ทั่วไป แต่เป็นสแตกการอนุมานที่ออกแบบมาเพื่อการควบคุมแบบเรียลไทม์

FlashRT ระบุตำแหน่งของตัวเองอย่างชัดเจนตั้งแต่ต้น README: มันคือเอนจินการอนุมานประสิทธิภาพสูงที่มุ่งเน้น “ปริมาณงาน AI แบบเรียลไทม์ที่มีขนาดแบตช์เล็กและความหน่วงต่ำ” ตำแหน่งนี้มีความสำคัญอย่างยิ่ง

การเพิ่มประสิทธิภาพการอนุมานแบบดั้งเดิมมักจะตั้งสมมติฐานสถานการณ์หนึ่งโดยปริยาย: เซิร์ฟเวอร์รับคำขออย่างต่อเนื่อง และเพิ่มปริมาณงานโดยรวมผ่านการประมวลผลแบบแบตช์ การจัดการ KV cache คิวจัดตารางงาน และกลยุทธ์ผู้เช่าหลายราย อย่างไรก็ตาม การควบคุม VLA ของหุ่นยนต์นั้นเหมือน “วงควบคุมทางอุตสาหกรรม” มากกว่า: การสังเกตจากกล้องและคำสั่งข้อความแต่ละครั้งจะต้องถูกแปลงเป็นลำดับการกระทำโดยเร็วที่สุด คอขวดที่นี่ไม่ได้มีแค่การคูณเมทริกซ์เท่านั้น แต่ครอบคลุมทั้งห่วงโซ่ ตั้งแต่การจัดตารางงานของ Python การกระจายโอเปอเรเตอร์ของเฟรมเวิร์ก รูปร่างไดนามิก การจัดสรรหน่วยความจำ การปรับเทียบเชิงปริมาณ ไปจนถึงการเลือก implement กลไก attention

โซลูชันของ FlashRT รวมถึง:

ใช้ CUDA kernel ที่เขียนด้วยมือเพื่อครอบคลุมเส้นทางที่ไวต่อหน่วยความจำ เช่น normalization, activation function, Rotary Position Embedding (RoPE), fused operator และ quantization
ใช้ cuBLASLt, CUTLASS และ FlashAttention-2 เพื่อจัดการกับ GEMM และกลไก attention ที่เน้นการคำนวณหนัก
ใช้ FP8 E4M3 และ NVFP4 เพื่อลดค่าใช้จ่ายด้านหน่วยความจำของน้ำหนักและแอคติเวชัน
ใช้ CUDA Graph เพื่อจับภาพการส่งต่อทั้งหมดเป็นกราฟแบบคงที่ หลังจากนั้น การอนุมานแต่ละครั้ง只需เล่นกราฟนั้นซ้ำ
ผ่านชั้นกระจายฮาร์ดแวร์ ซ่อนความแตกต่างของเป้าหมายต่างๆ เช่น Jetson Thor, RTX 5090, RTX 4090 ไว้ในแบ็กเอนด์

นี่ไม่ใช่แค่ “การเร่งความเร็วโมเดล PyTorch” แต่เป็นการออกแบบระบบอนุมานใหม่เป็นสถาปัตยกรรมที่ “ส่วนหน้ารับผิดชอบ I/O และน้ำหนัก ส่วนระดับล่างรับผิดชอบพอยน์เตอร์ที่เสถียรและกราฟแบบคงที่”

2.2 ข้อมูลประสิทธิภาพหลักเผยให้เห็นขอบเขตเป้าหมายของมัน

ผลลัพธ์ประสิทธิภาพที่เป็นตัวแทนที่แสดงใน README มีดังนี้:

Pi0.5 บน Jetson AGX Thor ประมาณ 44 มิลลิวินาที ประมาณ 23 Hz
Pi0.5 บน RTX 5090 ประมาณ 17.58 มิลลิวินาที ประมาณ 57 Hz
GROOT N1.6 บน RTX 5090 ประมาณ 13.08 มิลลิวินาที
Pi0-FAST บน RTX 5090 ประมาณ 2.39 มิลลิวินาที/โทเคน
Qwen3.6-27B NVFP4 บน RTX 5090 โดยทั่วไปประมาณ 100 โทเคน/วินาที สูงสุดประมาณ 129 โทเคน/วินาที และรองรับเส้นทางบริบท 256K

ความสำคัญของตัวเลขเหล่านี้ไม่ได้อยู่ที่ “ความเร็ว” เท่านั้น แต่ยังชี้ไปที่เป้าหมายทางวิศวกรรม: ในสถานการณ์ขนาดแบตช์เล็ก ไม่ต้องพึ่งพาการประมวลผลแบบแบตช์ที่ซับซ้อนเพื่อกระจายค่าใช้จ่าย แต่เป็นการกำจัดค่าใช้จ่ายเหล่านี้โดยตรง

สาม. สถาปัตยกรรมโดยรวม: แบ่งการส่งต่อโมเดลออกเป็นสัญญาสี่ชั้น

3.1 ชั้น API สาธารณะ: ผู้ใช้เพียงแค่สัมผัส load_model และ predict

จุดเริ่มต้นการเรียกใช้ของ FlashRT คือ flash_rt.load_model(...) ฟังก์ชันนี้รวบรวมพารามิเตอร์ต่างๆ มากมาย เช่น เส้นทาง checkpoint, ประเภทเฟรมเวิร์ก, การกำหนดค่าโมเดล, การเลือกฮาร์ดแวร์, กลยุทธ์ CUDA Graph และสวิตช์ FP4 ไว้ใน API ที่กระชับ

ที่มา: flash_rt/api.py  
def load_model(checkpoint, framework="torch", num_views=2, autotune=3,  
recalibrate=False, weight_cache=True, config="pi05", device=None,  
decode_cuda_graph=False, decode_graph_steps=80,  
max_decode_steps=256,  
hardware="auto",  
embodiment_tag=None,  
action_horizon=None,  
use_fp4=False,  
fp4_layers=None,  
use_awq=None,  
awq_alpha=0.5,  
use_p1_split_gu=None):  
"""โหลดโมเดล FlashRT

Args:
checkpoint: เส้นทางไปยังไดเรกทอรี checkpoint
            - torch: ไดเรกทอรี safetensors
            - jax: ไดเรกทอรี Orbax checkpoint
framework: "torch" หรือ "jax"
num_views: จำนวนมุมมองกล้อง (ค่าเริ่มต้น 2)
autotune: ความเข้มข้นของการปรับแต่งอัตโนมัติ CUDA Graph
config: ชื่อการกำหนดค่าโมเดล: "pi05", "pi0", "groot", "pi0fast"
hardware: การเลือกแบ็กเอนด์ GPU ``"auto"`` (ค่าเริ่มต้น) จะตรวจจับ
อุปกรณ์ CUDA ปัจจุบันผ่าน compute capability และเลือก
แบ็กเอนด์ที่ตรงกันที่สุด

เบื้องหลังโค้ดนี้คือปรัชญาการออกแบบ “เปลือกนอกที่เสถียร ภายในสามารถเปลี่ยนแทนกันได้” ผู้ใช้ไม่จำเป็นต้องสนใจว่า Pi0.5 บน Thor ใช้ CUTLASS FMHA หรือบน RTX ใช้ FlashAttention-2 และไม่จำเป็นต้องรู้ว่ารูปแบบน้ำหนักเป็น safetensors หรือ Orbax ผู้ใช้เพียงแค่แสดง “ฉันต้องการโหลดโมเดลไหน” ระบบจะเลือกเส้นทางที่เหมาะสมที่สุดโดยอัตโนมัติ

3.2 ชั้น Frontend: อินพุต น้ำหนัก และการปรับเทียบที่เกี่ยวข้องกับเฟรมเวิร์ก

FlashRT รองรับทั้งรูปแบบ PyTorch safetensors และ JAX Orbax อย่างสมบูรณ์ ความแตกต่างของเฟรมเวิร์กส่วนใหญ่อยู่ที่ส่วนหน้า: รวมถึงการอ่าน checkpoint, การเตรียม tokenizer, การประมวลผลภาพล่วงหน้า และการกระตุ้นการปรับเทียบ

ในเมธอด VLAModel.predict จะเห็นว่ามันจัดเรียงรายการภาพที่ผู้ใช้ป้อนเข้าเป็นพจนานุกรมการสังเกตก่อน จากนั้นจึงกระตุ้นการปรับเทียบตามข้อมูลจริงเมื่อจำเป็น และสุดท้ายเรียกใช้ฟังก์ชัน infer ของ pipeline ระดับล่าง:

ที่มา: flash_rt/api.py  

if self._needs_real_data_calibration:  
self._pipe.calibrate_with_real_data([obs])  
self._needs_real_data_calibration = False  

result = self._pipe.infer(obs)  
return result['actions']

กลยุทธ์ “การปรับเทียบแบบขี้เกียจ” (lazy calibration) นี้มีความชาญฉลาดทางวิศวกรรมอย่างยิ่ง: ชั้น API ยังคงความกระชับเพียงสามบรรทัด แต่การป้อนข้อมูลจริงครั้งแรกจะดำเนินการเตรียมการ เช่น การปรับขนาดแอคติเวชัน FP8 หลังจากนั้นก็จะเข้าสู่เส้นทางการเล่นกราฟซ้ำที่เสถียร

3.3 ชั้น Pipeline: การส่งต่อแบบพอยน์เตอร์ที่ไม่ขึ้นกับเฟรมเวิร์ก

เอกสารเทมเพลตโปรเจกต์อธิบายโมเดลทางความคิดของ FlashRT อย่างชัดเจน: forward() ของโมเดลดั้งเดิมจะถูกแปลเป็นชุดการเรียกใช้ kernel fvk.* โค้ด Python ชุดนี้จะถูกดำเนินการเพียงครั้งเดียวในระหว่างขั้นตอนการจับภาพ CUDA Graph เมื่ออนุมานจริง จะไม่มีการดำเนินการทีละบรรทัดอีกต่อไป แต่จะเล่นกราฟทั้งหมดซ้ำโดยตรง

ที่มา: flash_rt/frontends/torch/_template/README.md  
| โค้ดโมเดลของคุณ                           | ค่าที่เทียบเท่าใน FlashRT                                            |
| ----------------------------------------- | ------------------------------------------------------------- |
| `out = model.encoder(images, text)`       | `encoder_forward(ctx, fvk, bufs, weights, dims)`              |
| `F.rms_norm(x, weight)`                   | `fvk.rms_norm_fp16(x_ptr, w_ptr, out_ptr, S, D, eps, stream)`|
| `F.scaled_dot_product_attention(q, k, v)` | `attn.run("encoder", layer_idx, q_seq=S, stream=stream)`      |
| การเรียก `model(...)` ครั้งต่อๆ ไป             | `infer(obs)` → `cudaGraphLaunch(self._enc_ae_graph)`          |

**ข้อมูลเชิงลึกที่สำคัญ**: FlashRT แทนที่ `forward()` ของโมเดลของคุณด้วย  
CUDA Graph ที่ถูกจับภาพไว้อย่างสมบูรณ์

นี่คือลักษณะเฉพาะของคอมไพเลอร์ของ FlashRT: มันไม่ได้สร้าง IR ในความหมายดั้งเดิม และไม่ได้ส่งออก ONNX แต่เป็นการ “ลดระดับด้วยมือ” (hand-degrade) กระบวนการส่งต่อโมเดลเป็นกราฟโอเปอเรเตอร์แบบคงที่และโปรโตคอลพอยน์เตอร์ ในแง่หนึ่ง มันเหมือนกับแบ็กเอนด์คอมไพเลอร์แบบคงที่ที่ออกแบบมาโดยเฉพาะสำหรับตระกูลโมเดลจำนวนไม่กี่ตระกูล

3.4 ชั้น Hardware: attention และ GEMM จ่ายงานตามฮาร์ดแวร์

FlashRT ห่อหุ้มความแตกต่างของฮาร์ดแวร์ไว้ในไดเรกทอรี hardware/ คำอธิบายสถาปัตยกรรมใน README แสดงให้เห็นว่า Thor, RTX 5090, RTX 4090 จะถูกส่งต่อไปยังแบ็กเอนด์ที่แตกต่างกันโดยอัตโนมัติ โดยเฉพาะอย่างยิ่งในส่วนของ attention: Thor สามารถใช้เส้นทาง CUTLASS FMHA หรือ cuBLAS decomposition ในขณะที่ RTX ใช้ vendored FlashAttention-2

แกนหลักของนามธรรมชั้นนี้คือโปรโตคอล AttentionBackend

ที่มา: flash_rt/hardware/backend.py
class AttentionBackend(Protocol):
“””ผู้ให้บริการฮาร์ดแวร์สำหรับ attention ใน pipeline โมเดล”””

def sites(self) -> tuple[str, …]:
“””ส่งคืน tuple ของชื่อ site ที่แบ็กเอนด์นี้ถูกกำหนดค่าไว้”””

def get_slot_ptrs(self, site: str, layer_idx: int) -> dict[str, int]:
“””ส่งคืนพอยน์เตอร์อุปกรณ์ int ดิบสำหรับทุก slot ที่ (site, layer)

พอยน์เตอร์ที่ส่งคืนจะใช้ได้ตลอดอายุการใช้งานของแบ็กเอนด์และ
เสถียรตลอดการจับภาพและเล่นซ้ำของ CUDA Graph
“””
…

def run(
self,
site: str,
layer_idx: int,
q_seq: int,
*,
kv_seq: Optional[int] = None,
stream: int = 0,
state_nk: Optional[int] = None,
) -> int:
“””ดำเนินการ attention สำหรับหนึ่ง (site, layer) และส่งคืนพอยน์เตอร์เอาต์พุต”””

โปรโตคอลนี้แก้ปัญหาพื้นฐาน: pipeline โมเดลไม่จำเป็นต้องสนใจวิธีการ implement attention ที่เฉพาะเจาะจงเลย สิ่งที่มันต้องการคือชุดพอยน์เตอร์ slot Q/K/V ที่เสถียร และอินเทอร์เฟซการเรียกใช้ run() ที่เป็นหนึ่งเดียว แบ็กเอนด์สามารถจัดการ buffer ได้เอง หรือ pipeline จะเป็นผู้ถือ buffer ก็ได้ ตราบใดที่พอยน์เตอร์เหล่านี้ยังคงเสถียรในระหว่างการจับภาพและเล่นซ้ำของ CUDA Graph ก็สามารถรวมเข้าไปในกราฟแบบคงที่ได้อย่างราบรื่น

สี่. ทำไม CUDA Graph ถึงเป็นแกนหลักด้านประสิทธิภาพของ FlashRT

4.1 ในสถานการณ์ batch เล็ก การจัดตารางงานของ Python เองคือคอขวดด้านประสิทธิภาพ

เมื่อขนาดแบตช์ใหญ่พอ ค่าใช้จ่ายในการจัดตารางงานฝั่ง CPU สามารถถูกกลบด้วยการคำนวณของ GPU อย่างไรก็ตาม สถานการณ์การควบคุม VLA มักจะมี batch=1 หรือขนาดแบตช์ที่เล็กมาก ในกรณีนี้ การจัดตารางงานของแต่ละ PyTorch operation, พฤติกรรมของตัวจัดสรรหน่วยความจำ, การซิงค์ stream และลูป Python ล้วนกลายเป็นแหล่งที่มาของความหน่วงที่รับรู้ได้

FlashRT ใช้ ctypes เพื่อเรียกใช้ CUDA Runtime API โดยตรง ทำให้เกิด wrapper CUDA Graph ที่ไม่ขึ้นกับเฟรมเวิร์ก:

ที่มา: flash_rt/core/cuda_graph.py
class CUDAGraph:
“””CUDA Graph ที่ไม่ขึ้นกับเฟรมเวิร์กโดยใช้ CUDA Runtime API ดิบ”””

def __init__(self):
self._graph = ctypes.c_void_p()
self._graph_exec = ctypes.c_void_p()
self._captured = False

def begin_capture(self, stream: ctypes.c_void_p):
# cudaStreamCaptureModeRelaxed=2: จับภาพเฉพาะ operation บน STREAM นี้เท่านั้น
_check(_cudart.cudaStreamBeginCapture(stream, 2), “cudaStreamBeginCapture”)

def end_capture(self, stream: ctypes.c_void_p):
_check(_cudart.cudaStreamEndCapture(stream, ctypes.byref(self._graph)),
“cudaStreamEndCapture”)
_check(_cudart.cudaGraphInstantiate(
ctypes.byref(self._graph_exec), self._graph, 0),
“cudaGraphInstantiate”)
self._captured = True

def replay(self, stream: ctypes.c_void_p):
if not self._captured:
raise RuntimeError(“No graph captured”)
_check(_cudart.cudaGraphLaunch(self._graph_exec, stream), “cudaGraphLaunch”)

ที่นี่ใช้ cudaStreamCaptureModeRelaxed=2 โดยเจตนา โดยคอมเมนต์ระบุชัดเจนว่าจับภาพเฉพาะ operation บน stream ปัจจุบันเท่านั้น สิ่งนี้สำคัญอย่างยิ่งสำหรับการทำงานร่วมกับแบ็กเอนด์ JAX/XLA เนื่องจากโหมดการจับภาพทั่วโลกอาจขัดแย้งกับเธรดพื้นหลังของเฟรมเวิร์ก

4.2 ข้อกำหนดหลักในการจับภาพกราฟ: พอยน์เตอร์คงที่ ไม่มีการจัดสรรแบบไดนามิก รูปร่างคงที่

CUDA Graph ไม่ใช่เวทมนตร์ เหตุผลที่มันสามารถปรับปรุงประสิทธิภาพได้นั้นต้องแลกมาด้วย “ความเป็นคงที่”: ลำดับ kernel, พอยน์เตอร์ และความสัมพันธ์ในการจัดสรรหน่วยความจำที่บันทึกไว้ระหว่างการจับภาพ จะต้องยังคงใช้ได้ในระหว่างการเล่นซ้ำ ดังนั้น โค้ดจำนวนมากของ FlashRT จึงวนเวียนอยู่กับหลักการสำคัญข้อหนึ่ง: buffer กลางทั้งหมดถูกจัดสรรไว้ล่วงหน้า และเส้นทางการส่งต่อจะส่งผ่านเฉพาะพอยน์เตอร์ดิบเท่านั้น

CudaBuffer คือการแสดงออกที่เป็นรูปธรรมของแนวคิดการออกแบบนี้ในการ implement ระดับล่าง

# ที่มา: flash_rt/core/cuda_buffer.py
class CudaBuffer:
"""GPU buffer — managed or device memory."""

def __init__(self, nbytes: int, managed: bool = True):
self._ptr = ctypes.c_void_p()
self._managed = managed
if managed:
_check(_cudart.cudaMallocManaged(ctypes.byref(self._ptr), nbytes, 1),
"cudaMallocManaged")
else:
_check(_cudart.cudaMalloc(ctypes.byref(self._ptr), nbytes), "cudaMalloc")
self._nbytes = nbytes

@classmethod
def from_numpy(cls, arr: np.ndarray) -> 'CudaBuffer':
"""สร้าง device buffer, อัปโหลดผ่าน chunked H2D (Thor-safe ≤4MB/chunk)."""
arr = np.ascontiguousarray(arr)
buf = cls(arr.nbytes, managed=False)
chunk = 4 * 1024 * 1024  # 4MB chunks, Thor-safe
for off in range(0, arr.nbytes, chunk):
n = min(chunk, arr.nbytes - off)
_check(_cudart.cudaMemcpy(
ctypes.c_void_p(buf._ptr.value + off),
ctypes.c_void_p(arr.ctypes.data + off), n, 1), "H2D chunk")
return buf

คอมเมนต์ “Thor-safe ≤4MB/chunk” ในโค้ดน่าสนใจเป็นพิเศษ นี่ไม่ใช่ทฤษฎีนามธรรมในเอกสารวิชาการ แต่เป็นข้อจำกัดของฮาร์ดแวร์จากสถานการณ์การปรับใช้จริง ปรัชญาการออกแบบของ FlashRT ปรากฏชัดเจน ณ จุดนี้: เพื่อรับประกันความน่าเชื่อถือของระบบเรียลไทม์ การเพิ่มประสิทธิภาพหลายอย่างไม่ได้มุ่งหวัง “ความสวยงาม” แต่เป็น “ความแน่นอน” มากกว่า