BarraCUDA - คอมไพเลอร์ CUDA โอเพนซอร์สสำหรับ AMD GPU

• คอมไพเลอร์โอเพนซอร์สแบบสแตนด์อโลน ที่แปลง ซอร์ส CUDA C (.cu) เป็น โค้ดเครื่องสำหรับ AMD RDNA3 (GFX11) ได้โดยตรง
• สร้างไบนารี ELF .hsaco ผ่าน เล็กเซอร์, พาร์เซอร์ และตัวแทนกลาง (BIR) ของตัวเอง โดยไม่ต้องพึ่ง LLVM หรือชั้น HIP
• เขียนด้วย โค้ด C99 ราว 15,000 บรรทัด และสามารถบิลด์ได้ด้วยคำสั่ง make เพียงคำสั่งเดียว
• รองรับความสามารถหลักของ CUDA เช่น ตัวแปรบิลต์อินของเธรด, หน่วยความจำที่ใช้ร่วมกัน, การทำงานแบบอะตอมมิก, การทำงานระดับวาร์ป และ cooperative groups
• เผยแพร่ภายใต้ Apache 2.0 License และมีเป้าหมายจะขยายไปยังสถาปัตยกรรมอื่นเพิ่มเติมในอนาคต เช่น Tenstorrent, Intel Arc, RISC-V

---

ภาพรวมของ BarraCUDA

• BarraCUDA คือ คอมไพเลอร์ CUDA สำหรับ AMD GPU ที่แปลงไฟล์ .cu เป็น โค้ดเครื่อง GFX11
  • ผลลัพธ์อยู่ในรูปแบบ ไบนารี ELF .hsaco ที่สามารถรันบน AMD GPU ได้
  • ทำงานได้อย่างอิสระเต็มรูปแบบ โดยไม่ต้องพึ่ง LLVM
• โค้ดทั้งหมดมีขนาดประมาณ 15,000 บรรทัด เขียนด้วย C99 และบิลด์ได้ด้วย Makefile เดียว
• โปรเจ็กต์นี้พัฒนาโดย นักพัฒนาจากนิวซีแลนด์ ในลักษณะงานส่วนตัว

วิธีการทำงาน

• ประมวลผลไฟล์ .cu ที่ป้อนเข้ามาตามลำดับ พรีโปรเซส → วิเคราะห์คำศัพท์ → พาร์ส → วิเคราะห์เชิงความหมาย → สร้าง BIR → เลือกคำสั่ง → จัดสรรรีจิสเตอร์ → เข้ารหัสไบนารี → ส่งออก ELF
• BIR (BarraCUDA IR) เป็นตัวแทนภายในแบบ SSA ที่ออกแบบให้ไม่ผูกกับสถาปัตยกรรมใดสถาปัตยกรรมหนึ่ง
• การเข้ารหัสทั้งหมดตรวจสอบด้วย llvm-objdump แล้ว และมี ข้อผิดพลาดในการดีโค้ด 0 รายการ

ความสามารถที่รองรับ

• ไวยากรณ์หลักของ CUDA: __global__, __device__, __host__, threadIdx, blockIdx เป็นต้น
• ฟีเจอร์ของ CUDA: หน่วยความจำ __shared__, __syncthreads(), การทำงานแบบอะตอมมิก, warp shuffle/vote, ชนิดข้อมูลเวกเตอร์, ความแม่นยำแบบ half, cooperative groups เป็นต้น
• ความสามารถของคอมไพเลอร์: C preprocessor ที่สมบูรณ์, การกู้คืนจากข้อผิดพลาด, การติดตามตำแหน่งในซอร์ส, รองรับการส่งค่าโครงสร้าง

สิ่งที่ยังไม่รองรับ

• การใช้ unsigned แบบเดี่ยว, ตัวดำเนินการกำหนดค่าแบบผสม (+=, -= เป็นต้น), const, หน่วยความจำ __constant__, อาร์เรย์ shared แบบ 2D, texture·surface, dynamic parallelism ยังไม่ได้ถูกติดตั้งใช้งาน
• ยังไม่รองรับหลาย translation units และการสร้างโค้ดฝั่งโฮสต์

การทดสอบและโรดแมป

• ตรวจสอบแล้วด้วย ไฟล์ทดสอบ 14 ไฟล์, เคอร์เนลมากกว่า 35 ตัว, และ โค้ดเครื่องราว 27KB
• เป้าหมายระยะสั้น: เติมความสมบูรณ์ของไวยากรณ์และเพิ่มความเข้ากันได้กับไฟล์ .cu สำหรับการใช้งานจริง
• เป้าหมายระยะกลาง: ปรับปรุง instruction scheduling, การจัดสรรรีจิสเตอร์, constant folding, loop-invariant code motion และการเพิ่มประสิทธิภาพอื่น ๆ
• เป้าหมายระยะยาว: เพิ่มแบ็กเอนด์ใหม่สำหรับ Tenstorrent, Intel Arc, RISC-V Vector Extension เป็นต้น

ไลเซนส์

• Apache 2.0 License

[GN⁺]

ความคิดเห็นจาก Hacker News

• รู้สึกประทับใจกับอารมณ์ขันแบบนิวซีแลนด์ที่เขียนไว้ใน README ของโปรเจ็กต์
  การที่โปรเจ็กต์นี้ทำ instruction encoding ขึ้นมาเองโดยไม่พึ่ง LLVM ถือว่าสดใหม่มาก
  มันแสดงให้เห็นว่าการจะเริ่มโปรเจ็กต์แบบนี้ต้องมีความรู้ระดับล่างมาก
  ฝั่ง AMD มักถูกมองว่าไม่รองรับ CUDA จึงกลายเป็นข้ออ้างให้ NVIDIA ผูกขาด และความพยายามแบบนี้น่าจะช่วยสร้างสมดุลให้ตลาดได้

  • มีคนทักว่าที่ต้นฉบับพูดถึงไม่ใช่ LLM แต่เป็น LLVM
  • คิดว่าโปรเจ็กต์นี้ก็น่าจะมีการมีส่วนร่วมที่เกี่ยวกับ AIอยู่บ้าง แต่ถ้าใช้อย่างเหมาะสม AI ก็เป็นเครื่องมือที่ยอดเยี่ยมได้
  • พอเห็นพูดถึง “อารมณ์ขันแบบโอเชียเนีย” ก็ทำให้นึกถึง แอนิเมชัน Beached Whale
  • คอมเมนต์ในซอร์สโค้ดที่ว่า “/* 80 keywords walk into a sorted bar */” เล่นคำได้เฉียบมาก (ลิงก์ lexer.c)
  • เห็นคำว่า “AI slope” แล้วมีคนต่อมุกว่าตามหลักวิทยาศาสตร์มันควรเรียกว่า gradient descent

• แปลกใจที่ external issue แรกถูกเปิดโดย geohot (ลิงก์ issue)
  อยากเห็นคนแบบนี้ร่วมมือกันเพื่อทำลายการผูกขาดตลาด GPUของ NVIDIA

• การรันเวิร์กโหลด AI inference บน GPU ของ NVIDIA มีต้นทุนสูง
  โปรเจ็กต์แบบนี้สำคัญต่อการสร้างทางเลือกที่สตาร์ทอัพพอจะจ่ายไหว
  อยากรู้ว่าประสิทธิภาพของโอเปอเรชันอย่าง conv2d หรือ attention จะเป็นอย่างไร

• “# It’s C99. It builds with gcc. There are no dependencies.”
  ความเรียบง่ายที่จบได้ด้วย make บรรทัดเดียวช่างงดงามจริง ๆ

  • ชอบแนวทางที่เรียบง่ายและตรงไปตรงมาแบบนี้มาก

• พอเห็นการใช้ตัวพิมพ์ใหญ่ในชื่อโพสต์ ก็เพิ่งเข้าใจว่าทำไม GPU farm ของบริษัทถึงชื่อว่า “barracuda” มันขำดีทีเดียว

• ถ้านักพัฒนาที่มีแพสชันทำในสิ่งที่ AMD ทำไม่ได้สำเร็จ ก็คงเป็นเรื่องที่ทั้งขำและเศร้า

  • คิดว่า AMD ไม่รองรับ CUDA ไม่ใช่เพราะ “ทำไม่ได้” แต่เป็นการเลือกเชิงกลยุทธ์
    เพราะถ้ารองรับ CUDA ก็จะยิ่งไปเสริมความแข็งแกร่งให้ ecosystem ของ NVIDIA
    ถ้าอยากได้ตัวแทน CUDA จริง ๆ ZLUDA อาจใช้งานได้จริงมากกว่า
  • โอเพนซอร์สมักไม่ต้องคอยกังวลสายตาผู้ถือหุ้น จึงบางครั้งให้ผลลัพธ์ที่ดีกว่าเชิงพาณิชย์
    แต่ก็น่าเสียดายที่พอโตขึ้นแล้วมักสุดท้ายถูกบริษัทใหญ่ซื้อไปและหายไป
    เหมือนกรณีของ Linus กับ git ที่แสดงให้เห็นว่าถ้ามีทั้งความมุ่งมั่นและความรู้ ก็สามารถพังกำแพงได้
  • หลายกรณี AMD ไม่ใช่ “ทำไม่ได้” แต่เป็น “ไม่ทำ”
    เช่น การไม่รองรับ FSR4 บนการ์ดรุ่นเก่าอย่างเป็นทางการก็เป็นตัวอย่างหนึ่ง
  • ปิดท้ายด้วยมุกว่า “เรามี HIP นะ”

• สงสัยว่าจะรองรับสถาปัตยกรรม AMD รุ่นเก่า (GFX1010 เป็นต้น) ได้ไหม

  • คิดว่ายากแต่เป็นไปได้แน่นอน เมื่อก่อน ROCm แย่มากจนต้องแพตช์เอง แต่ตอนนี้ดีขึ้นเยอะแล้ว
  • ฉันก็ใช้ RX5700m อยู่เลยกำลังทำงานเรื่องการรองรับ AMD รุ่นเก่าด้วยตัวเอง
    กำลังอ่านเอกสาร ISA แล้วปรับ binary encoding อยู่
    แต่พื้นที่นี้เป็นงานที่LLM ไม่ค่อยเก่ง จึงต้องเข้าใจและแก้เอง

• CUDA รองรับ C++ เลยมีคนสงสัยว่าการเดินหน้าแบบ pure C โดยไม่มี Clang/LLVM จะไม่จำกัดเกินไปหรือ

  • ในทางปฏิบัติตอนนี้กำลัง parse แค่**บางส่วนของฟีเจอร์ C++**ที่ CUDA ใช้ ตัวคอมไพเลอร์เองเขียนด้วย C99
  • จากการทดสอบ ดูเหมือนว่าจะรองรับฟีเจอร์ C++ บางอย่างอย่าง template ได้ด้วย
  • การรองรับ AMD GX11 ของ LLVM มีจำกัดมาก จึงดีกว่าหากไม่ต้องพึ่งมัน
  • นักพัฒนาของจริงเขียนโค้ดโดยไม่พึ่ง AI
    ทุกวันนี้โอเพนซอร์สเต็มไปด้วย PR ที่ AI สร้างขึ้น จึงน่าประทับใจที่โปรเจ็กต์นี้หลีกเลี่ยงการพึ่งพาแบบนั้น

• คิดว่า AMD ควรสนับสนุนอย่างเป็นทางการให้กับโปรเจ็กต์แบบนี้
  โลกจำเป็นต้องหลุดพ้นจากการผูกขาดของ NVIDIA

• มีคนสงสัยว่า “OpenCL ตายไปแล้วหรือยัง?”
  ไม่ค่อยแน่ใจ แต่ถึงอย่างนั้นโปรเจ็กต์นี้ก็ยังน่าประทับใจ

  • CUDA กินส่วนแบ่งตลาดไปมากกว่าอย่างชัดเจน และ OpenCL มีฐานที่อ่อนแอ
    เมื่อก่อน Apple ก็เคยรองรับ แต่ตอนนี้เหมือนจะไม่แล้ว
    คล้ายความสัมพันธ์ระหว่าง Unix กับ Windows คือในเชิงเทคนิคอาจดี แต่ตลาดกลับเทไปด้านเดียว