BarraCUDA - AMD GPU के लिए ओपन सोर्स CUDA कंपाइलर

• CUDA C source (.cu) को AMD RDNA3(GFX11) के लिए machine code में सीधे बदलने वाला standalone open source compiler
• LLVM या HIP layer के बिना अपने lexer, parser, intermediate representation (BIR) के जरिए ELF .hsaco binary बनाता है
• लगभग 15,000 lines के C99 code में लिखा गया है, और एक ही make command से build किया जा सकता है
• CUDA के thread built-in variables, shared memory, atomic operations, warp operations, cooperative groups जैसी प्रमुख सुविधाओं को support करता है
• Apache 2.0 license के तहत जारी किया गया है, और आगे Tenstorrent, Intel Arc, RISC-V जैसी अतिरिक्त architectures तक विस्तार का लक्ष्य है

---

BarraCUDA का परिचय

• BarraCUDA AMD GPU के लिए CUDA compiler है, जो .cu files को GFX11 machine code में बदलता है
  • इसका output AMD GPU पर चलने योग्य ELF .hsaco binary के रूप में होता है
  • यह LLVM dependency के बिना पूरी तरह standalone तरीके से काम करता है
• पूरा codebase C99 में लिखी गई लगभग 15,000 lines का है, और एक single Makefile से build किया जा सकता है
• इस project को New Zealand के एक developer ने व्यक्तिगत रूप से विकसित किया है

यह कैसे काम करता है

• दी गई .cu file को preprocessing → lexical analysis → parsing → semantic analysis → BIR generation → instruction selection → register allocation → binary encoding → ELF output क्रम में process किया जाता है
• BIR(BarraCUDA IR), SSA form वाला internal representation है, जिसे architecture-independent तरीके से design किया गया है
• सभी encodings को llvm-objdump के जरिए verify किया गया है और decode errors 0 हैं

समर्थित सुविधाएँ

• CUDA core syntax: __global__, __device__, __host__, threadIdx, blockIdx आदि
• CUDA features: __shared__ memory, __syncthreads(), atomic operations, warp shuffle/vote, vector types, half precision, cooperative groups आदि
• Compiler features: पूरा C preprocessor, error recovery, source location tracking, struct value passing support

अभी समर्थित नहीं

• unsigned का standalone उपयोग, compound assignment operators (+=, -= आदि), const, __constant__ memory, 2D shared arrays, texture·surface, dynamic parallelism आदि अभी implement नहीं किए गए हैं
• multiple translation units और host code generation supported नहीं हैं

टेस्ट और रोडमैप

• 14 test files, 35 से अधिक kernels, और लगभग 27KB machine code के साथ verify किया गया है
• अल्पकालिक लक्ष्य: syntax को पूरा करना और वास्तविक उपयोग वाली .cu files के साथ compatibility मजबूत करना
• मध्यम अवधि के लक्ष्य: instruction scheduling, register allocation improvements, constant folding, loop-invariant code motion जैसी optimizations
• दीर्घकालिक लक्ष्य: Tenstorrent, Intel Arc, RISC-V Vector Extension जैसे नए backends जोड़ना

लाइसेंस

• Apache 2.0 license

[GN⁺]

Hacker News की राय

• प्रोजेक्ट के README में लिखा New Zealand-स्टाइल humor काफ़ी प्रभावशाली लगा
  LLVM पर निर्भर हुए बिना खुद का instruction encoding लागू करना ताज़गीभरा लगा
  यह दिखाता है कि ऐसा प्रोजेक्ट शुरू करने के लिए कितनी गहरी low-level knowledge चाहिए
  AMD पक्ष में CUDA सपोर्ट न होना अक्सर NVIDIA के एकाधिकार का बहाना बन जाता है, और ऐसा प्रयास बाज़ार में संतुलन लाने में मददगार लग सकता है

  • बताया गया कि मूल पाठ में LLM नहीं बल्कि LLVM कहा गया था
  • लगा कि इस प्रोजेक्ट में भी निश्चित रूप से AI से जुड़ा contribution शामिल होगा। लेकिन सही तरीके से इस्तेमाल किया जाए तो AI एक बेहतरीन tool हो सकता है
  • “Oceania-स्टाइल humor” की बात सुनकर Beached Whale animation याद आ गया
  • source code की टिप्पणी “/* 80 keywords walk into a sorted bar */” काफ़ी चतुर लगी (lexer.c लिंक)
  • “AI slope” अभिव्यक्ति देखकर मज़ाक में जोड़ा कि वैज्ञानिक रूप से इसे gradient descent कहा जाता है

• यह देखकर हैरानी हुई कि पहला बाहरी issue geohot ने उठाया था (issue लिंक)
  ऐसे लोगों को मिलकर NVIDIA के GPU बाज़ार एकाधिकार को तोड़ते देखना अच्छा लगेगा

• NVIDIA GPU पर AI inference workload चलाना काफ़ी महँगा पड़ता है
  ऐसे प्रोजेक्ट startup कंपनियों के लिए वहन योग्य विकल्प बनाने में महत्वपूर्ण हैं
  conv2d या attention जैसे operations में performance कैसी होगी, यह जानने की जिज्ञासा है

• “# It’s C99. It builds with gcc. There are no dependencies.”
  make की एक पंक्ति में सब हो जाने वाली सादगी सच में खूबसूरत है

  • यह सरल और सहज तरीका बहुत पसंद आया

• पोस्ट के शीर्षक में बड़े अक्षरों का इस्तेमाल देखकर अब जाकर समझ आया कि कंपनी के GPU farm का नाम “barracuda” क्यों है। काफ़ी मज़ेदार लगा

• अगर जोशीले developers वह कर दिखाएँ जो AMD नहीं कर सकी, तो वह मज़ेदार भी होगा और दुखद भी

  • लगा कि AMD ने CUDA को सपोर्ट न करने का कारण “न कर पाना” नहीं बल्कि रणनीतिक फैसला था
    क्योंकि CUDA सपोर्ट करने से उल्टा NVIDIA ecosystem और मज़बूत हो जाता
    अगर CUDA का विकल्प चाहिए, तो ZLUDA ज़्यादा व्यावहारिक हो सकता है
  • open source को shareholders की चिंता नहीं करनी पड़ती, इसलिए वह कई बार commercial products से बेहतर नतीजे देता है
    लेकिन अफ़सोस यह है कि बड़ा होने पर अक्सर किसी बड़ी कंपनी द्वारा खरीदा जाकर ग़ायब हो जाता है
    Linus और git की तरह, अगर इच्छा और ज्ञान हो तो दीवारें तोड़ी जा सकती हैं
  • AMD के कई मामले ऐसे हैं जहाँ उसने “नहीं किया”, न कि “नहीं कर सकी”
    उदाहरण के लिए FSR4 का पुराने cards पर आधिकारिक सपोर्ट न देना भी ऐसा ही मामला है
  • अंत में “हमारे पास HIP है” वाले मज़ाक से बात खत्म की गई

• जिज्ञासा हुई कि क्या पुराने AMD architectures (GFX1010 आदि) तक सपोर्ट बढ़ाया जा सकेगा

  • मुश्किल ज़रूर है, लेकिन काफ़ी संभव लगता है। पहले ROCm बहुत खराब था इसलिए खुद patch करना पड़ता था, लेकिन अब काफ़ी बेहतर है
  • मैं भी RX5700m इस्तेमाल करता हूँ, इसलिए पुराने AMD सपोर्ट पर खुद काम कर रहा हूँ
    ISA documents पढ़ते हुए binary encoding को समायोजित कर रहा हूँ
    लेकिन यह ऐसा क्षेत्र है जहाँ LLM ज़्यादा अच्छे नहीं होते, इसलिए खुद समझकर ठीक करना पड़ता है

• CUDA में C++ सपोर्ट है, इसलिए यह सवाल उठा कि Clang/LLVM के बिना pure C वाला रास्ता कहीं सीमित तो नहीं होगा

  • दरअसल अभी CUDA में इस्तेमाल होने वाली C++ features के सिर्फ़ कुछ हिस्सों को parse किया जा रहा है। compiler खुद C99 में लिखा गया है
  • tests देखकर लगा कि templates जैसे कुछ C++ features भी संभाले जा सकते हैं
  • LLVM का AMD GX11 सपोर्ट बहुत सीमित है, इसलिए उस पर निर्भर न रहना बेहतर है
  • असली developers AI पर निर्भर हुए बिना code लिखते हैं
    आजकल open source में AI-जनित PR की बाढ़ है, इसलिए इस प्रोजेक्ट का ऐसी निर्भरता से बचना प्रभावशाली लगा

• लगा कि AMD को ऐसे प्रोजेक्ट्स को आधिकारिक sponsorship देनी चाहिए
  दुनिया को NVIDIA के एकाधिकार से बाहर निकलने की ज़रूरत है

• यह सवाल उठा कि “क्या OpenCL अब खत्म हो चुका है?”
  पक्का नहीं, लेकिन फिर भी यह प्रोजेक्ट प्रभावशाली है

  • CUDA ने कहीं ज़्यादा market share हासिल कर लिया है, जबकि OpenCL की नींव कमज़ोर है
    पहले Apple भी इसे सपोर्ट करता था, लेकिन अब शायद नहीं
    कुछ-कुछ Unix और Windows के रिश्ते जैसा है—तकनीकी रूप से अच्छा होने के बावजूद बाज़ार एक ही तरफ झुक गया