Tiny GPU: GPU ขั้นต่ำที่สร้างด้วย Verilog

 (github.com/adam-maj)
2 คะแนน โดย GN⁺ 2024-04-27 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • tiny-gpu คือการอิมพลีเมนต์ GPU ขั้นต่ำบนพื้นฐานของ Verilog สำหรับเรียนรู้จากระดับล่างสุดว่า GPU ทำงานอย่างไรในระดับฮาร์ดแวร์ โดยเน้นหลักการร่วมกันของ GPGPU และตัวเร่ง ML มากกว่าฮาร์ดแวร์เฉพาะงานกราฟิก
  • การอิมพลีเมนต์ประกอบด้วย ไฟล์ Verilog ที่มีเอกสารกำกับน้อยกว่า 15 ไฟล์, เอกสารสถาปัตยกรรมและ ISA, เคอร์เนลบวกและคูณเมทริกซ์, รวมถึงการรองรับการจำลองเคอร์เนลและการติดตามการทำงาน
  • GPU รันได้ทีละ เคอร์เนลเดียว โดยโหลดหน่วยความจำโปรแกรมและหน่วยความจำข้อมูล ตั้งค่า thread_count แล้วส่งสัญญาณ start เพื่อเริ่มเคอร์เนล
  • เพื่อให้เรียบง่าย แต่ละคอร์จะประมวลผล block ได้ครั้งละหนึ่งชุด และแต่ละเธรดมี ALU, LSU, PC และ register file ของตัวเอง แต่สมมติว่าเธรดทั้งหมดจะกลับมารวมที่ PC เดียวกันหลังทุกคำสั่ง
  • ฟีเจอร์ส่วนใหญ่ของ GPU สมัยใหม่ เช่น แคชหลายชั้น, shared memory, memory coalescing, pipelining, warp scheduling, branch divergence และ barrier ถูกตัดออกเพื่อให้ความสำคัญกับโครงสร้างสำหรับการเรียนรู้

ปัญหาที่ tiny-gpu พยายามแก้

  • ฝั่ง CPU มีสื่อให้เรียนรู้ตั้งแต่สถาปัตยกรรมไปจนถึงสัญญาณควบคุมอยู่มาก แต่รายละเอียดทางเทคนิคระดับต่ำของ GPU สมัยใหม่ส่วนใหญ่ยังเป็นกรรมสิทธิ์เพราะตลาดมีการแข่งขันสูง
  • แม้จะมีสื่อการเขียนโปรแกรม GPU มากมาย แต่แทบไม่มีสื่อที่ช่วยให้เรียนรู้ว่า GPU ทำงานอย่างไรใน ระดับฮาร์ดแวร์
  • อิมพลีเมนต์ GPU แบบโอเพนซอร์สอย่าง Miaow และ VeriGPU มุ่งให้ทำงานได้ครบและใช้งานจริง จึงมีโครงสร้างซับซ้อน
  • tiny-gpu ตัดความซับซ้อนจำนวนมากของการ์ดจอระดับโปรดักชันออก แล้วโฟกัสกับองค์ประกอบแกนหลักที่ตัวเร่งฮาร์ดแวร์สมัยใหม่มีร่วมกัน
    • องค์ประกอบสำคัญของสถาปัตยกรรม GPU
    • วิธีที่โมเดลการเขียนโปรแกรมแบบ SIMD ถูกอิมพลีเมนต์ในฮาร์ดแวร์
    • วิธีที่ GPU จัดการกับแบนด์วิดท์หน่วยความจำที่จำกัด

สถาปัตยกรรมโดยรวม

  • tiny-gpu ถูกออกแบบให้รันได้ ทีละเคอร์เนลเดียว
  • ขั้นตอนการรันเคอร์เนลมีดังนี้
    • โหลดโค้ดเคอร์เนลเข้า global program memory
    • โหลดข้อมูลที่ต้องใช้เข้า data memory
    • ระบุจำนวนเธรดที่จะรันใน device control register
    • ตั้งสัญญาณ start เป็น high เพื่อเริ่มรันเคอร์เนล
  • GPU ประกอบด้วยยูนิตต่อไปนี้
    • device control register
    • dispatcher
    • compute core จำนวนแปรผัน
    • memory controller สำหรับ data memory และ program memory
    • cache

การรันเคอร์เนลและการกระจายเธรด

  • device control register ทำหน้าที่เก็บเมทาดาทาของการรันเคอร์เนล และใน tiny-gpu จะเก็บเพียง thread_count ซึ่งเป็นจำนวนเธรดทั้งหมดที่จะรัน
  • dispatcher จะกระจายเธรดไปยัง compute core หลายตัวเมื่อเริ่มเคอร์เนล
    • จัดกลุ่มเธรดที่สามารถรันแบบขนานได้เป็น block
    • ส่ง block ไปให้คอร์ที่พร้อมใช้งานประมวลผล
    • แจ้งว่าการรันเคอร์เนลเสร็จสิ้นเมื่อประมวลผลทุก block ครบแล้ว
  • คอร์แบบย่อส่วนนี้ประมวลผลได้ครั้งละหนึ่ง block
  • แต่ละเธรดมี ALU, LSU, PC และ register file เฉพาะตัว
  • การจัดการการรันคำสั่งของเธรดบนทรัพยากรเหล่านี้เป็นหนึ่งในปัญหาที่ยากของ GPU

โครงสร้างหน่วยความจำและคอนโทรลเลอร์

  • GPU ถูกสร้างให้เชื่อมต่อกับ global memory ภายนอก และเพื่อความเรียบง่ายจึงแยก data memory ออกจาก program memory
  • สเปกของ data memory
    • addressable แบบ 8 บิต
    • มีทั้งหมด 256 แถว
    • ข้อมูลขนาด 8 บิต
    • แต่ละแถวเก็บค่าได้ต่ำกว่า 256
  • สเปกของ program memory
    • addressable แบบ 8 บิต
    • มีทั้งหมด 256 แถว
    • ข้อมูลขนาด 16 บิต
    • แต่ละคำสั่งมีขนาด 16 บิตตาม ISA
  • memory controller จะติดตามคำขอหน่วยความจำจากคอร์ จำกัดคำขอให้สอดคล้องกับแบนด์วิดท์ของหน่วยความจำภายนอกจริง และส่งต่อคำตอบกลับไปยังทรัพยากรที่ถูกต้อง
  • memory controller แต่ละตัวมีจำนวนช่องสัญญาณคงที่ตามแบนด์วิดท์ของ global memory
  • cache ยังอยู่ระหว่างพัฒนา โดยจะเก็บข้อมูลที่ดึงมาจากหน่วยความจำภายนอกไว้ใน SRAM ภายในอุปกรณ์ เพื่อให้คำขอครั้งถัดไปดึงได้เร็วขึ้น และเปิดแบนด์วิดท์หน่วยความจำไว้สำหรับข้อมูลใหม่

องค์ประกอบภายในคอร์

  • แต่ละคอร์มี scheduler ตัวเดียวสำหรับจัดการการรันของเธรด
  • scheduler ของ tiny-gpu จะรันคำสั่งของหนึ่ง block ไปจนจบก่อนจึงค่อยรับ block ใหม่ และรันคำสั่งของทุกเธรดตามลำดับที่ซิงก์กัน
  • scheduler ที่ก้าวหน้ากว่านี้สามารถเพิ่มอัตราการใช้ทรัพยากรด้วย pipelining และ warp scheduling
  • ข้อจำกัดหลักของ scheduler คือ latency ที่เกิดขึ้นเมื่อโหลดและบันทึกข้อมูลจาก global memory
    • คำสั่งส่วนใหญ่รันแบบ synchronous ได้
    • งาน load-store อย่าง LDR และ STR เป็นแบบ asynchronous จึงต้องจัดโครงสร้างการรันคำสั่งรอบช่วงรอที่ยาว
  • Fetcher จะดึงคำสั่งของ program counter ปัจจุบันจาก program memory แบบ asynchronous
  • Decoder จะแปลงคำสั่งที่ดึงมาเป็นสัญญาณควบคุมสำหรับการรันของเธรด
  • register file ของแต่ละเธรดใช้เก็บข้อมูลระหว่างคำนวณและทำให้เกิดแพตเทิร์น SIMD
    • register แบบอ่านอย่างเดียวมี %blockIdx, %blockDim, %threadIdx
    • เคอร์เนลสามารถรันกับข้อมูลที่ต่างกันตาม local thread ID
  • ALU ของแต่ละเธรดรองรับคำสั่งเลขคณิต ADD, SUB, MUL, DIV
  • CMP จะให้ผลว่าค่าผลต่างของสองรีจิสเตอร์เป็นลบ ศูนย์ หรือบวก และเก็บผลไว้ในรีจิสเตอร์ NZP ของยูนิต PC
  • LSU ของแต่ละเธรดเข้าถึง global data memory และจัดการ LDR, STR รวมถึงเวลารอหน่วยความจำแบบ asynchronous
  • PC ของแต่ละเธรดกำหนดคำสั่งถัดไปที่จะรัน
    • โดยปกติจะเพิ่มขึ้นทีละ 1 ต่อหนึ่งคำสั่ง
    • BRnzp จะกระโดดไปยังแถวของ program memory ที่กำหนด หากเงื่อนไขในรีจิสเตอร์ NZP ที่ตั้งโดย CMP ก่อนหน้าตรงกัน
    • ลูปและเงื่อนไขถูกอิมพลีเมนต์ด้วยวิธีนี้
  • เพื่อความเรียบง่าย tiny-gpu สมมติว่าเธรดทั้งหมดจะกลับมารวมที่ PC เดียวกันหลังทุกคำสั่ง
  • ใน GPU จริง แต่ละเธรดสามารถแตกแขนงไปยัง PC ที่ต่างกันได้ และเมื่อนั้นจะเกิด branch divergence ซึ่งทำให้กลุ่มเธรดที่เคยถูกประมวลผลร่วมกันแยกออกเป็นหลาย execution flow

ISA

  • tiny-gpu อิมพลีเมนต์ ISA 11 คำสั่ง เพื่อรันเคอร์เนลเรียบง่ายสำหรับพิสูจน์แนวคิด เช่น การบวกเมทริกซ์และการคูณเมทริกซ์
  • คำสั่งที่รองรับ
    • BRnzp: กระโดดไปยังแถวอื่นของ program memory หากเงื่อนไข NZP ตรงกัน
    • CMP: เปรียบเทียบค่าในสองรีจิสเตอร์และบันทึกผลลงรีจิสเตอร์ NZP
    • ADD, SUB, MUL, DIV: การคำนวณเลขคณิตพื้นฐานสำหรับ tensor math
    • LDR: โหลดข้อมูลจาก global memory
    • STR: บันทึกข้อมูลลง global memory
    • CONST: โหลดค่าคงที่เข้ารีจิสเตอร์
    • RET: ส่งสัญญาณจบการรันของเธรดปัจจุบัน
  • แต่ละรีจิสเตอร์ถูกระบุด้วย 4 บิต จึงมีทั้งหมด 16 รีจิสเตอร์
    • R0 ถึง R12 จำนวน 13 ตัวเป็นรีจิสเตอร์อิสระที่อ่านและเขียนได้
    • 3 ตัวสุดท้ายเป็น special register แบบอ่านอย่างเดียวที่ให้ค่า %blockIdx, %blockDim, %threadIdx ซึ่งจำเป็นต่อ SIMD

ลำดับการรัน

  • เมื่อแต่ละคอร์รันคำสั่ง จะใช้ลำดับการควบคุมในขั้นต่อไปนี้
    • FETCH: ดึงคำสั่งถัดไปของ PC ปัจจุบัน
    • DECODE: ถอดรหัสคำสั่งเป็นสัญญาณควบคุม
    • REQUEST: ขอข้อมูลจาก global memory หากเป็น LDR หรือ STR
    • WAIT: รอการตอบกลับจาก global memory หากจำเป็น
    • EXECUTE: ทำการคำนวณกับข้อมูล
    • UPDATE: อัปเดต register file และรีจิสเตอร์ NZP
  • ลำดับการควบคุมนี้ถูกออกแบบมาเพื่อความเรียบง่ายและเข้าใจได้ง่าย
  • ในการอิมพลีเมนต์จริง อาจยุบรวมบางขั้นตอนเพื่อปรับเวลาในการประมวลผลให้เหมาะสม หรือใช้ pipelining เพื่อจัดการการรันหลายคำสั่งบนทรัพยากรของคอร์
  • แต่ละเธรดจะคำนวณตาม execution path เดียวกันโดยอาศัยข้อมูลใน register file เฉพาะของตน
  • แผนภาพคล้ายกับของ CPU แต่ต่างกันตรงที่มี %blockIdx, %blockDim, %threadIdx อยู่ในรีจิสเตอร์แบบอ่านอย่างเดียวเพื่อเปิดใช้ความสามารถแบบ SIMD

ตัวอย่างเคอร์เนล

  • มีการเขียนเคอร์เนลบวกเมทริกซ์และคูณเมทริกซ์เพื่อพิสูจน์แนวคิดของ ISA
  • ไฟล์ทดสอบในรีโพซิทอรีสามารถจำลองเคอร์เนลเหล่านี้บน GPU ได้ครบถ้วน และสร้างสถานะของ data memory พร้อม execution trace ทั้งหมด

  • การบวกเมทริกซ์

    • matadd.asm ใช้บวกเมทริกซ์ขนาด 1 x 8 สองชุด
    • การบวกทีละองค์ประกอบทั้ง 8 จะทำในเธรดแยกกัน
    • ใช้รีจิสเตอร์ %blockIdx, %blockDim, %threadIdx เพื่อแสดงการเขียนโปรแกรมแบบ SIMD
    • ใช้คำสั่ง LDR และ STR จึงครอบคลุมการจัดการหน่วยความจำแบบ asynchronous

  • การคูณเมทริกซ์

    • matmul.asm ใช้คูณเมทริกซ์ขนาด 2 x 2 สองชุด
    • คำนวณ dot product ของแถวและคอลัมน์ที่เกี่ยวข้องทีละองค์ประกอบ
    • ใช้ CMP และ BRnzp เพื่อแสดงการแตกแขนงภายในเธรด
    • เนื่องจากทุกแขนงกลับมารวมกันอีกครั้ง จึงทำงานได้ในอิมพลีเมนต์ tiny-gpu ปัจจุบัน

การจำลอง

  • หากต้องการรันการจำลองเคอร์เนล ต้องมี iverilog และ cocotb
  • ขั้นตอนเตรียมพร้อม
    • ติดตั้ง Verilog compiler และ cocotb ด้วย brew install icarus-verilog และ pip3 install cocotb
    • ดาวน์โหลดเวอร์ชันล่าสุดของ sv2v แตกไฟล์ และเพิ่มไบนารีลงใน $PATH
    • รัน mkdir build ที่รูทของรีโพซิทอรี
  • การจำลองเคอร์เนลรันได้ด้วย make test_matadd และ make test_matmul
  • ผลการรันจะถูกเขียนเป็นไฟล์ล็อกใน test/logs
    • สถานะเริ่มต้นของ data memory
    • execution trace ทั้งหมดของเคอร์เนล
    • สถานะสุดท้ายของ data memory
  • ตอนต้นของแต่ละไฟล์ล็อกจะแสดงเมทริกซ์อินพุต และส่วนท้ายใน data memory สุดท้ายจะแสดงเมทริกซ์ผลลัพธ์
  • execution trace มีสถานะการรันของทุกเธรดในทุกคอร์สำหรับแต่ละ cycle
    • คำสั่งปัจจุบัน
    • PC
    • ค่ารีจิสเตอร์
    • ข้อมูลสถานะ

ความสามารถขั้นสูงของ GPU ที่ตั้งใจไม่ใส่มา

  • tiny-gpu ตัดองค์ประกอบเพิ่มประสิทธิภาพและความสามารถส่วนใหญ่ของ GPU สมัยใหม่ออกเพื่อให้เรียบง่าย

  • แคชหลายชั้นและ shared memory

    • GPU สมัยใหม่ใช้ลำดับชั้นของแคชหลายระดับเพื่อลดการเข้าถึง global memory
    • tiny-gpu อิมพลีเมนต์เพียงชั้นแคชเดียวที่เก็บข้อมูลล่าสุดไว้ระหว่างทรัพยากรที่ร้องขอกับ memory controller
    • แคชหลายชั้นช่วยเก็บข้อมูลที่ใช้บ่อยให้ใกล้ตำแหน่งใช้งานมากขึ้นเพื่อลดเวลาโหลด
    • GPU ยังอาจใช้ shared memory เพื่อให้เธรดใน block เดียวกันแลกเปลี่ยนผลลัพธ์ที่ใช้ร่วมกันได้

  • memory coalescing

    • ระหว่างการรันแบบขนาน หลายเธรดมักเข้าถึงแอดเดรสต่อเนื่องกัน เช่น องค์ประกอบที่อยู่ติดกันของเมทริกซ์
    • memory coalescing จะวิเคราะห์คำขอหน่วยความจำในคิวและรวมคำขอที่อยู่ติดกันให้เป็นทรานแซกชันเดียว
    • เป้าหมายคือเพื่อลดเวลาในการอ้างแอดเดรสและประมวลผลคำขอร่วมกัน

  • pipelining

    • คอร์ของ tiny-gpu จะเริ่มคำสั่งถัดไปได้ก็ต่อเมื่อการรันคำสั่งหนึ่งของชุดเธรดก่อนหน้าสิ้นสุดแล้ว
    • GPU สมัยใหม่สามารถสตรีมการรันคำสั่งตามลำดับหลายคำสั่งได้ พร้อมยังคงรับประกันลำดับสำหรับคำสั่งที่มีการพึ่งพากัน
    • ช่วยเพิ่มอัตราการใช้ทรัพยากรไม่ให้คอร์ว่างในช่วงอย่างเช่นการรอคำขอหน่วยความจำแบบ asynchronous

  • warp scheduling

    • แบ่ง block ออกเป็น warp ซึ่งเป็นชุดเธรดที่สามารถรันร่วมกันได้
    • เมื่อ warp หนึ่งกำลังรอ ก็จะรันคำสั่งของ warp อื่น ทำให้คอร์เดียวประมวลผลหลาย warp พร้อมกันได้
    • คล้ายกับ pipelining แต่เป็นการจัดการคำสั่งจากคนละเธรด

  • branch divergence

    • tiny-gpu สมมติว่าเธรดทั้งหมดในชุดเดียวกันจะมี PC เดียวกันหลังแต่ละคำสั่ง
    • ในความเป็นจริง แต่ละเธรดสามารถแตกแขนงไปยังบรรทัดที่ต่างกันตามข้อมูลได้
    • เธรดที่มี PC ต่างกันจะถูกแยกเป็น execution flow คนละชุด และยังต้องจัดการจุดที่กลับมารวมกันอีกครั้ง

  • การซิงก์และ barrier

    • GPU สมัยใหม่สามารถตั้ง barrier เพื่อให้กลุ่มเธรดใน block เดียวกันรอจนทุกเธรดมาถึงจุดที่กำหนด
    • มีประโยชน์เมื่อเธรดต้องแลกเปลี่ยนข้อมูลร่วมกัน และต้องการรับประกันว่าการประมวลผลข้อมูลเสร็จสมบูรณ์แล้ว

งานถัดไป

  • รายการปรับปรุงในอนาคตมีดังนี้
    • เพิ่ม instruction cache แบบง่าย
    • สร้าง adapter เพื่อให้ใช้ GPU บน Tiny Tapeout 7 ได้
    • เพิ่ม branch divergence ขั้นพื้นฐาน
    • เพิ่ม memory coalescing ขั้นพื้นฐาน
    • เพิ่ม pipelining ขั้นพื้นฐาน
    • ปรับปรุง control flow และการใช้รีจิสเตอร์เพื่อให้ cycle time ดีขึ้น
    • เขียนเคอร์เนลกราฟิกพื้นฐานหรือเพิ่มฮาร์ดแวร์กราฟิกแบบง่ายเพื่อสาธิตความสามารถด้านกราฟิก
  • ผู้ใช้ที่ต้องการช่วยพัฒนารีโพซิทอรีสามารถมีส่วนร่วมผ่าน PR ได้

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 2024-04-27
ความคิดเห็นจาก Hacker News
  • ตลาด GPU มีการแข่งขันสูงมาก ทำให้ รายละเอียดทางเทคนิคระดับต่ำ ของสถาปัตยกรรมสมัยใหม่ส่วนใหญ่ยังคงไม่เปิดเผย
    ข้อยกเว้นคือ Intel ที่เผยแพร่เอกสารเทคนิคเกี่ยวกับ GPU ไว้มากมาย: https://kiwitree.net/~lina/intel-gfx-docs/prm/
    คู่มือ i810/815 ก็หาได้ออนไลน์ และถ้าไม่นับช่องว่างแปลก ๆ ในช่วงก่อน 965 ที่ขาด 855/910/915/945 ไป การจัดทำเอกสารก็ถือว่าค่อนข้างต่อเนื่อง
    • AMD ก็เผยแพร่เอกสารไว้ไม่น้อยเช่นกัน: https://www.amd.com/en/developer/browse-by-resource-type/documentation.html
      รวมถึง เอกสารสถาปัตยกรรมชุดคำสั่ง ของผลิตภัณฑ์ปัจจุบันและในอดีตด้วย แต่ดูเหมือนจะเป็นเอกสารสำหรับผู้พัฒนาการใช้งานจริงมากกว่าคำอธิบายระดับสูงสำหรับนักเล่นที่สนใจ
    • ไดรเวอร์ Linux ของ Intel ก็มีคุณภาพดีและอยู่ใน mainline
      อยากให้ทุกบริษัททำตามแนวทางนี้
    • แม้จะเป็นข้อมูลจากปี 2018 แต่ก็ยังเกี่ยวข้องอยู่บ้าง: The Thirty Million Line Problem - Casey Muratori
  • เป็นโปรเจกต์ที่เจ๋งมากจริง ๆ และดีใจที่ได้เห็นโปรเจกต์ฮาร์ดแวร์แบบนี้ดำเนินแบบเปิดเผย
    แต่ผมมองว่านี่ใกล้เคียงกับ โปรเซสเซอร์ร่วม SIMD มากกว่า
    ถ้าจะเรียกว่า GPU อย่างน้อยก็ควรต้องมีเอาต์พุตแสดงผลบางรูปแบบ
    ผมรู้ว่าช่วงหลังคำนี้ถูกใช้หลวมขึ้นมาก เพราะ Nvidia และบริษัทอื่น ๆ ขายสถาปัตยกรรมกราฟิกเวอร์ชันสำหรับเซิร์ฟเวอร์โดยเฉพาะในชื่อ GPU ด้วย แต่ในงานออกแบบ GPU ส่วนของกราฟิกก็ยังคงเป็นสัดส่วนความซับซ้อนที่มากอยู่ดี
    • ถ้ามันประมวลผลกราฟิก ผมคิดว่าก็ถือเป็น GPU ได้แม้ไม่มีเอาต์พุต
      GPU ที่ไม่ส่งภาพออกมาก็ยังมีประโยชน์
      ที่ทำงานของผมมีเวิร์กสเตชันประมาณ 75 เครื่องที่ใส่ Quadro ระดับกลางไว้ การ์ดมีแต่ mini-DisplayPort แต่บริษัทซื้อให้แค่สาย HDMI ทุกเครื่องเลยต่อเข้ากราฟิกออนบอร์ดกันหมด
      ถึงอย่างนั้นการ์ดพวกนั้นก็ยังเร่งซอฟต์แวร์และประมวลผลกราฟิก เพียงแค่ไม่ได้ส่งภาพขึ้นจอเท่านั้น
  • ดีมาก ผมสนับสนุนงาน โอเพนคอร์ GPU อย่างเต็มที่
    มีตัวอย่างอื่นด้วย: https://github.com/jbush001/NyuziProcessor
    • ถ้ามี การ implement CUDA แบบขั้นต่ำ สำหรับโปรเซสเซอร์โอเพนคอร์พวกนี้สักตัวก็คงดี
      ถ้าจะให้ TSMC หรือฟาวน์ดรีอื่นผลิตโปรเซสเซอร์แบบนี้ได้คุ้มทุน ต้องมีปริมาณประมาณเท่าไร?
  • เป็นโปรเจกต์ที่ยอดเยี่ยมจริง ๆ
    ผมอยากลอง FPGA แต่พูดตรง ๆ แค่จะเริ่มตรงไหนยังจับทางยาก และทั้งสาขาก็ดูน่าเกรงขามมาก
    เป้าหมายสุดท้ายคือทำการ์ดเร่งความเร็วสำหรับ LLM ซึ่งแม้จะเป็นเป้าหมายที่ตั้งขึ้นค่อนข้างตามใจ แต่ก็น่าจะมีส่วนทับซ้อนกับโปรเจกต์นี้เยอะ และอาจต่างกันแค่ส่วน offloading หน่วยความจำเพื่อโหลดโมเดลที่ใหญ่ขึ้นเท่านั้น
    • ต้องปรับกรอบความคิดในหัวใหม่
      การเริ่มต้นกับ FPGA ต้องแบ่งออกเป็นเทคโนโลยีย่อยหลายส่วน และต้องปรับความคาดหวังด้วย
      เราไม่คาดหวังให้วิศวกรซอฟต์แวร์สร้างคอมพิวเตอร์ทั้งเครื่องจากหลักการพื้นฐาน เขียนสถาปัตยกรรมชุดคำสั่ง เข้าใจภาษาเครื่อง แปลงมันเป็นแอสเซมบลี และพัฒนาภาษาโปรแกรมเพื่อสร้างแอปด้วยโค้ด Python ตั้งแต่แรก
      ควรเริ่มจากด้านบนแล้วค่อยไล่ลงไปตามสแต็ก
      ถ้าโฟกัสที่การซ่อนความซับซ้อนและสร้างระบบจาก IP ที่ทำไว้แล้ว การออกแบบ FPGA ก็จะค่อนข้างง่าย
      โดยทั่วไปจะแนะนำเครื่องมืออย่าง MATLAB เพราะถ้าใช้ DevKit ที่มี reference design ก็สามารถสร้างแอปเริ่มต้นด้วย HDL Coder ได้
      ไม่อย่างนั้นจะต้องแบกรับภาระมหาศาลในการเรียนรู้สถาปัตยกรรมคอมพิวติ้งดิจิทัล, Verilog, timing, transceiver/I/O, การวางแผนขา, Quartus/Vivado, simulation/verification, ระบบ embedded ฯลฯ
      สรุปคือให้เริ่มจาก การออกแบบระดับระบบ เรียนรู้วิธีนำ IP แบบ plug-and-play มาเชื่อมต่อกันที่ระดับบนสุด แล้วลองใส่โมดูลนั้นเข้าไปใน reference design ที่ทำไว้แล้ว
      จากนั้นค่อย ๆ ลอกชั้นออกเพื่อเผยให้เห็นความซับซ้อนด้านล่าง
    • ผมก็อยู่ในสถานการณ์เดียวกัน และแผนคือแบบนี้
      1. อ่าน Digital Design and Computer Architecture ของ Harris, Harris (2022), Elsevier: https://doi.org/10.1016/c2019-0-00213-0
      2. ทำตามคอร์ส RVFpga ของผู้เขียน เพื่อสร้าง RISC-V CPU จริงบน FPGA: https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • แนะนำเส้นทางแบบนี้
      1. โคลนรีโพซิทอรีเพื่อการศึกษา https://github.com/yuri-panchul/basics-graphics-music เป็นชุดแล็บง่าย ๆ สำหรับคนที่เริ่มเรียน Verilog ตั้งแต่ศูนย์ เขียนโดย Yuri Panchul ซึ่งเคยทำงานที่ Imagination เพื่อพัฒนา GPU
      2. หา FPGA board สักรุ่นจากหลายสิบรุ่นที่รองรับ พร้อมอุปกรณ์เสริมอย่างปุ่มกดและ LED
      3. ติดตั้ง Yosys และเครื่องมือที่เกี่ยวข้อง
      4. เริ่มจาก lab01 DeMorgan แล้วทำแล็บในรีโพซิทอรีให้มากที่สุดเท่าที่ทำได้
        สามารถทำแล็บควบคู่ไปกับการอ่าน Harris&Harris ได้
        เมื่อทำแล็บและอ่านหนังสือจบ ก็ถึงเวลาจะเริ่มโปรเจกต์ของตัวเอง
        อนึ่ง HackerMojo มีการพบปะรายสัปดาห์ด้วย และแม้ไม่ได้อยู่ใน Valley ก็เข้าร่วมผ่าน Zoom ได้
    • ไม่รู้ว่าคุณอยู่ระดับไหนแล้ว แต่แหล่งข้อมูลเหล่านี้ช่วยให้ผมเข้าใจ ตรรกะดิจิทัล และสถาปัตยกรรม CPU/GPU ได้ดีขึ้น
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
  1. https://hdlbits.01xz.net/wiki/Main_Page
  • ถ้าอยากเร่งความเร็ว LLM ก่อนอื่นต้องเข้าใจ สถาปัตยกรรม
    เริ่มจากตรงนั้นได้เลย
    จริง ๆ แล้วฮาร์ดแวร์ก็เป็นส่วนที่ง่าย และในแง่การผลิตก็เป็นส่วนที่ยากด้วย
  • มีเหตุผลอะไรไหมที่ในบล็อก always แบบลำดับตรงนี้ถึงใช้ตัวดำเนินการ การกำหนดค่าแบบ non-blocking ปนกับ การกำหนดค่าแบบ blocking?
    • นั่นดูเหมือนตัวแปรโลคัล
    • ถ้าไม่ได้ยึดติดมากเกินไปกับการให้ผลลัพธ์ของ simulation กับ synthesis ตรงกัน ก็ทำแบบนั้นได้
  • นานมาแล้วเคยทำอะไรคล้าย ๆ กันด้วย VHDL
    เคยมีเว็บชื่อ opencores ที่รวมโปรเจกต์ HDL โอเพนซอร์สหลาย ๆ อันไว้
    สงสัยว่าสมัยนี้มี simulator HDL แบบกระจายขนาดใหญ่ระดับ HPC ที่ดี ๆ บ้างไหม
    การใช้ GPU สมัยใหม่กับ simulation ระดับ RTL ดูสมเหตุสมผล
    • ไม่ใช่ “เคยมี” แต่ยังมีอยู่: https://opencores.org/projects?language=VHDL
      หรือว่าไม่ใช่เว็บเดียวกัน แต่เป็นเว็บอื่นที่คล้ายกัน?
  • ALU จะ implement คำสั่ง DIV ตรง ๆ ในระดับฮาร์ดแวร์เลยเหรอ?
    ในที่อย่าง CUDA core สมัยใหม่ ปกติแล้วมีการหารเป็นคำสั่งจริง ๆ ไหม หรือโดยทั่วไป emulate ด้วยซอฟต์แวร์?
    วงจรหารในฮาร์ดแวร์จริงกินพื้นที่มหาศาล เลยไม่คิดว่ามันจะอยู่ใน ALU ของ GPU
    ใน Verilog การเขียนแค่บรรทัดเดียวว่า DIV: begin alu_out_reg <= rs / rt; end นั้นง่ายมาก แต่บรรทัดเดียวนั้นกินซิลิคอนไปเยอะ
    ถ้าแค่ simulate Verilog อย่างเดียว ก็อาจมองไม่เห็นข้อเท็จจริงนี้
    • นี่เป็นแค่ โปรเจกต์สำหรับคนเรียน Verilog
      โปรเจกต์หยุดอยู่ที่ simulation และถ้าจะทำเป็นฮาร์ดแวร์จริงยังต้องทำงานอีกมาก
  • แถมยังเป็น “GPU” ที่ไม่มีฟังก์ชันกราฟิกอีก
    ส่วนตัวคิดว่าของแบบนี้ควรเรียกด้วยชื่ออื่น
    • คำถามแรกคือ ทำไมตั้งแต่แรก CPU กับ GPU ถึงถูกแยกออกจากกัน
      ช่องว่างระหว่างสองอย่างนี้แคบลงเรื่อย ๆ และทั้งสองฝ่ายต่างก็เพิ่มความสามารถของอีกฝ่ายเข้าไป แต่ก็ยังมีความแตกต่างค่อนข้างมาก
      ผมคิดว่ามันเกี่ยวข้องกับ กฎของ Amdahl [0]
      ในความหมายนี้ CPU อาจเรียกได้ว่าเป็นโปรเซสเซอร์ที่ปรับให้เหมาะกับ latency ส่วน GPU เป็นโปรเซสเซอร์ที่ปรับให้เหมาะกับ throughput
      ถ้าเจาะจงกว่านั้น [1] ก็อาจเรียก CPU ว่าเป็นโปรเซสเซอร์แบบ data dependency ที่ยาวและลึก และเรียก GPU ว่าเป็นโปรเซสเซอร์แบบ data dependency ที่กว้างและแบน
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • เรียกว่า TPU หรือ หน่วยประมวลผลเทนเซอร์ ก็ได้
      เทนเซอร์ก็คืออาร์เรย์ n มิติ
      สามารถวางซอฟต์แวร์หรือเฟิร์มแวร์ทับลงไปเพื่อให้ทำงานเหมือน GPU ได้
    • เคยคิดจะเริ่มโปรเจกต์ทำ ‘display adapter’ แต่ติดตั้งแต่ก่อนเริ่ม เพราะยังหาวิธีทำความเข้าใจโปรโตคอลสื่อสารระหว่าง GOP driver ของ UEFI กับ display adapter ไม่ได้
      พยายามปะติดปะต่อจากซอร์สของ EDK2 แล้ว แต่ไม่ชัดเจนว่าส่วนไหนเฉพาะกับ QEMU มากน้อยแค่ไหน
    • เรียกว่า MPU หรือ หน่วยประมวลผลเมทริกซ์ ก็ได้
    • คำที่ดูเหมือนกำลังตั้งหลักได้คือ AIA หรือ AI accelerator
  • การที่ tiny-gpu สมมติว่าเธรดทั้งหมดจะ “บรรจบ” กลับมาที่ program counter เดียวกันหลังแต่ละคำสั่ง เป็นการลดทอนที่ใสซื่อเกินไปมาก
    ใน GPU จริง เธรดแต่ละตัวสามารถ branch ไปยัง PC ที่ต่างกันได้ และจะเกิด branch divergence ที่กลุ่มเธรดซึ่งตอนแรกประมวลผลร่วมกันแยกออกไป execute ต่างหาก
    ก่อนจะทำ GPU เป็นซิลิคอน น่าจะลองเขียนโปรแกรม GPU ก่อน
    แถมจะเรียกว่า SIMD ก็ยังรู้สึกไม่ค่อยตรง
    คนนี้คือคนเดียวกับที่เมื่อก่อนเอาวงจรของคนอื่นมาต่อ ๆ กันให้ LED กะพริบ แล้วบอกว่าตัวเองทำ CPU นั่นแหละ
    • ข้อแรกนี่เทียบได้กับการเรียก __syncthreads() ทุกครั้งที่ execute เลยไม่ใช่หรือ?