HipKittens: तेज़ और शक्तिशाली AMD कर्नेल

 (hazyresearch.stanford.edu)
3 पॉइंट द्वारा GN⁺ 2025-11-16 | 1 टिप्पणियां | WhatsApp पर शेयर करें
  • HipKittens एक ऐसा प्रोजेक्ट है जो AMD GPU के लिए हाई-परफॉर्मेंस कर्नेल और C++-आधारित प्रोग्रामिंग primitives प्रदान करता है, ताकि AI कंप्यूटेशन की दक्षता बढ़ाई जा सके
  • मौजूदा AMD इकोसिस्टम के AITER, PyTorch, Triton, TileLang, Composable Kernel आदि ने अस्थिर परफॉर्मेंस और अपरिपक्व सपोर्ट जैसी सीमाएँ दिखाई हैं
  • HipKittens, tile abstraction को केंद्र में रखकर, NVIDIA और AMD के बीच एक साझा interface बनाए रखते हुए हार्डवेयर-विशिष्ट implementation को अलग करता है
  • लगभग 500 लाइनों से कम कोड में लिखे गए कर्नेल ने AMD के मौजूदा हस्तलिखित assembly कर्नेल से भी तेज़ परफॉर्मेंस हासिल की
  • यह AI कंप्यूटेशन को मल्टी-सिलिकॉन वातावरण तक विस्तार देने के लिए एक व्यावहारिक आधार प्रस्तुत करता है और ओपन हार्डवेयर इकोसिस्टम की ओर बदलाव की संभावना दिखाता है

AMD GPU इकोसिस्टम की वर्तमान स्थिति और समस्याएँ

  • अब तक AI कंप्यूटेशन का विकास मुख्यतः एकल हार्डवेयर वेंडर-केंद्रित रहा है, लेकिन AMD GPU नवीनतम पीढ़ी में शीर्ष-स्तरीय कंप्यूट परफॉर्मेंस और मेमोरी बैंडविड्थ प्रदान करते हैं
    • उदाहरण: AMD MI355X OAM में BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, मेमोरी 288GB, बैंडविड्थ 8.0TB/s
  • लेकिन परिपक्व software stack की कमी के कारण वास्तविक AI workflow में इस परफॉर्मेंस का पूरा उपयोग नहीं हो पाता
  • AMD के प्रमुख software components में AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK) शामिल हैं
    • AITER और PyTorch के SDPA Llama GQA backward कर्नेल, SoTA की तुलना में क्रमशः केवल 30% और 24% परफॉर्मेंस ही हासिल करते हैं
    • Mojo, bank conflict के कारण, लगभग 50% परफॉर्मेंस तक ही सीमित रहता है
    • TileLang केवल CDNA3 तक सपोर्ट करता है, और मुख्य फीचर्स की कमी तथा CK dependency के कारण जटिलता बढ़ती है
    • Triton को register lifetime tracking और memory access optimization में कठिनाई होती है
  • परिणामस्वरूप, सबसे उच्च-परफॉर्मेंस AMD कर्नेल अब भी विशेषज्ञों को सीधे assembly में लिखने पड़ते हैं, जो scalability और maintenance पर गंभीर बाधाएँ डालता है

CUDA-केंद्रित इकोसिस्टम से तुलना

  • AI समुदाय में इसे CUDA moat यानी CUDA इकोसिस्टम की प्रवेश-बाधा के रूप में देखा जाता है
  • अतीत में NVIDIA कर्नेल development के लिए भी लो-लेवल CUDA/CUTLASS पर आधारित वर्षों की मेहनत की आवश्यकता होती थी
  • हाल के वर्षों में DSL (domain-specific language) और AI-सहायक टूल्स के विकास से NVIDIA कर्नेल development सरल हुआ है
    • उदाहरण: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon आदि
  • इसी प्रवृत्ति के आधार पर AMD के लिए भी नई programming primitives की आवश्यकता की खोज की गई

HipKittens के डिज़ाइन सिद्धांत

  • HipKittens, ThunderKittens (NVIDIA) और ThunderMittens (Apple Silicon) के बाद AMD के लिए विकसित किया गया संस्करण है
  • मुख्य अवधारणा: tile abstraction
    1. tile abstraction की सार्वभौमिकता – NVIDIA में प्रभावी रहा tile-आधारित ऑपरेशन मॉडल AMD पर भी स्वाभाविक रूप से लागू होता है
    2. backend implementation की architecture-specific प्रकृति – memory access pattern और register scheduling हार्डवेयर के अनुसार अलग डिज़ाइन किए जाते हैं
    3. scheduling strategy की अनुकूलन क्षमता – AMD के CDNA3·CDNA4 में wave-आधारित scheduling अक्षम है, लेकिन tile-स्तरीय scheduling फिर भी सरलता और परफॉर्मेंस बनाए रखती है
  • interface (tile और operations) तथा implementation (hardware mapping) को अलग करके, यह विभिन्न GPU architectures पर समान रूप से लागू होने वाला मॉडल प्रस्तुत करता है

परफॉर्मेंस परिणाम

  • Attention Forward कर्नेल: लगभग 500 लाइनों के कोड में लिखा गया, और AITER के assembly कर्नेल से बेहतर परफॉर्मेंस हासिल की
    • विभिन्न head dimension (D) और sequence length (N) में, Causal/Non-Causal दोनों सेटिंग्स में बेहतर परिणाम
  • GEMM कर्नेल: 100 लाइनों से कम के मुख्य loop से बना है, और BF16 व FP8 दोनों ऑपरेशंस में सर्वोच्च परफॉर्मेंस हासिल करता है
  • Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm कर्नेल में भी मौजूदा सर्वोत्तम परफॉर्मेंस की तुलना में सुधार देखा गया
  • सभी कर्नेल पठनीयता और संशोधन में आसानी बनाए रखते हुए भी हस्तलिखित assembly-स्तर की परफॉर्मेंस प्रदान करते हैं

मल्टी-सिलिकॉन AI की ओर विस्तार

  • AI की पूरी क्षमता को साकार करने के लिए विविध और खुले हार्डवेयर के उपयोग की आवश्यकता है
  • HipKittens का लक्ष्य AMD GPU को AI डेवलपर्स के लिए वास्तव में सुलभ प्लेटफ़ॉर्म बनाना है
  • AMD के open source software stack के साथ मिलकर, यह मल्टी-सिलिकॉन आधारित AI इकोसिस्टम की ओर संक्रमण की संभावना दिखाता है
  • अगली पोस्ट में HipKittens की तकनीकी आंतरिक संरचना पर चर्चा की जाएगी (part two का संकेत)

संबंधित पढ़ाई

1 टिप्पणियां

 
GN⁺ 2025-11-16
Hacker News की राय

  • हमने AMD के साथ कॉन्ट्रैक्ट किया है और Llama 405B को MI350X पर MLPerf के लिए train कर रहे हैं
    AMD की स्थिति वाकई बेहतर हो रही है। अगर आपके पास AMD GPU है, तो pytorch.org पर Linux+ROCm चुनकर PyTorch इंस्टॉल करने की सलाह दूँगा। 3 साल पहले तक हालात निराशाजनक थे, लेकिन अब ज़्यादातर mainline features ठीक काम करते हैं। मैंने MI300X पर nanochat चलाया और वह सीधे काम कर गया। MI350X भी इसी तरह stable है
    बेशक यह अभी भी NVIDIA से पीछे है, और software ecosystem, compiler, और driver में बहुत निवेश की ज़रूरत है। लेकिन 2 साल पहले की निराशाजनक स्थिति की तुलना में अब उम्मीद दिख रही है।
    अगर आप AMD के LLVM backend की कमज़ोरियाँ देखना चाहते हैं, तो HipKittens के code की CUDA Kittens से तुलना करना अच्छा रहेगा।
    training के लिए NVIDIA और Google पहले स्थान पर हैं, AMD दूसरे पर है, और बाकी लगभग मौजूद ही नहीं हैं। Intel और Tenstorrent अभी बहुत पीछे हैं, Huawei में example segfault से crash हो गया, Groq ने chip बेचना छोड़ दिया, Cerebras कहीं उपलब्ध ही नहीं है। Trainium में एक instance पाने में 5 दिन लगे, इसलिए मेरी दिलचस्पी खत्म हो गई

    • मैं जानना चाहता हूँ कि Tinygrad को इस तरह के memory optimization या warp-specific features को व्यक्त या explore करने में अभी कितना समय लगेगा। यह भी जानना चाहता हूँ कि tinygrad में ऐसे features जोड़ना कितना जटिल होगा
    • मैं जानना चाहता हूँ कि consumer hardware (non-MI) पर भी ROCm + PyTorch चलाने के लिए proprietary kernel driver चाहिए या नहीं
    • “Cerebras कहीं उपलब्ध नहीं है” यह बात तो उल्टा जीत का संकेत लगती है
    • मैं पिछले 2 साल से AMD NeoCloud का CEO हूँ। AMD को इस तरह turnaround करते हुए सीधे देखना बहुत अच्छा लग रहा है।
      शुरुआती setup अभी भी थोड़ा rough है, लेकिन पहले की तुलना में बहुत बेहतर हुआ है। उदाहरण के लिए, एक महीने पहले तक nanochat ठीक से नहीं चल रहा था, लेकिन अब चलता है। महत्वपूर्ण बात यह है कि अब लोग AMD ecosystem में रुचि ले रहे हैं।
      AI को hardware diversity चाहिए। सारे AI hardware और software पर एक ही कंपनी का एकाधिकार शेयरधारकों के लिए अच्छा हो सकता है, लेकिन तकनीकी प्रगति के लिए हानिकारक है

  • मुझे NVIDIA का market cap समझ नहीं आता। अभी Transformer जैसे कुछ ही algorithm जीत गए हैं, और data ज़्यादा महत्वपूर्ण हो गया है। पहले की तरह विविध HPC code की ज़रूरत नहीं रही, इसलिए अब competitors को बस algorithm के एक सीमित set को optimize करना है।
    अगर vLLM और Transformer को efficiently चलाया जा सके, तो बहुत बड़ा बाज़ार खुलता है। ऐसे में AMD या Huawei भी आसानी से पकड़ बना सकते हैं, तो NVIDIA की असल moat क्या है? क्या सिर्फ InfiniBand काफ़ी है?

    • हाँ, NVIDIA की moat धीरे-धीरे कमज़ोर हो रही है। MS, Google, Amazon, Apple जैसी बड़ी कंपनियाँ सभी NVIDIA dependency से बचने के लिए अपने chip बना रही हैं।
      data center GPU में NVIDIA अभी भी मज़बूत है, लेकिन Google बड़े पैमाने पर TPU इस्तेमाल कर रहा है और OpenAI भी AMD hardware ऑर्डर कर रहा है।
      CUDA ecosystem अभी भी महत्वपूर्ण है, लेकिन उससे बाहर निकलने की कोशिश तेज़ हो चुकी है। AMD, Intel, Qualcomm आदि भी इस प्रतिस्पर्धा में उतर चुके हैं। HipKittens neutral software ecosystem की ओर एक महत्वपूर्ण कदम लगता है
    • कुछ algorithms को implement करने का सबसे आसान तरीका general-purpose compute hardware बनाना है। hardware का development cycle लंबा होता है, इसलिए यह approach स्थिर रहती है
    • InfiniBand की जगह UEC ले रहा है। inference के लिए InfiniBand की ज़रूरत नहीं होती। इसलिए inference market में कोई moat नहीं है। AMD या Google TPU का इस्तेमाल करना समझदारी है
    • NVIDIA का असली हथियार इसका विशाल CUDA library ecosystem है। लगभग हर क्षेत्र के लिए code मौजूद है
    • Transformer कोई एक तकनीक नहीं है। इसे implement करने के बहुत अलग-अलग तरीके हैं। इसलिए vLLM या TRL उतने सरल नहीं हैं

  • लगता है AMD को ऐसे project को fund करना चाहिए था, लेकिन मेरी राय में शायद उसने फिर मौका गंवा दिया। GPU और AI के मामले में वह हमेशा ऐसा ही करता रहा है

    • AMD product launch के लिए जितना न्यूनतम software चाहिए, बस उतना ही निवेश करता है। performance testing system तक कमज़ोर है, और regression bugs सीधे customers तक पहुँच जाते हैं।
      HipKittens सुधार है, लेकिन AMD के अंदर kernel performance track करने की क्षमता की कमी है। DevOps भारत की TCS को outsource किया गया था, और वहाँ के लोगों को हालात की ठीक समझ नहीं थी।
      जिन teams में अच्छे leaders हैं, वे अपने स्तर पर shadow IT teams चलाती हैं। ROCm के पास ऐसी team नहीं थी, और आखिरकार बड़े cloud customers के विरोध के बाद सुधार शुरू हुआ।
      talent hire करने पर भी market से कम salary ऑफ़र की जाती है, क्योंकि तुलना Qualcomm या Walmart स्तर से की जाती है।
      पिछले 4 साल में bonus कभी पूरा नहीं मिला
    • “AMD मौका नहीं छोड़ता, वह छूटा हुआ मौका छोड़ता है” वाली बात एकदम सही लगती है। Instinct hardware असल में NVidia से प्रतिस्पर्धा करने लायक है, लेकिन software support बेहद खराब था।
      पहले Radeon VII की तरह support बंद कर देना, या ecosystem को बार-बार बदल देना, इसकी वजह से maturity नहीं आ पाई। CUDA compatibility की कमी और निवेश की कमी के कारण ज़्यादातर organizations ने बस NVIDIA चुन लिया।
      फिर भी हाल के वर्षों में ROCm और Instinct ecosystem में लगातार निवेश हुआ है, इसलिए चीज़ें धीरे-धीरे बेहतर हो रही हैं। लेकिन networking में NVIDIA अभी भी बहुत आगे है
    • performance charts देखें तो AMD अब तक NVIDIA के स्तर पर हो सकता था। लेकिन software की कमी से वह हार गया। chip design ज़्यादा कठिन काम है, लेकिन असली समस्या software को न समझना थी
    • पहले कुछ लोगों ने ROCm-optimized kernels contribute किए थे, लेकिन AMD ने PR बंद कर दिए और merge नहीं किया। यह सच में समझ से बाहर है
    • अब funding मिल रही है और project सही तरह चल रहा है

  • मैं जानना चाहता हूँ कि ThunderKittens के ऊपर बने कोई project हैं या नहीं।
    अगर यह उसका HIP port है, तो साधारण CUDA port की तुलना में इसकी व्यावहारिक उपयोगिता कहीं ज़्यादा लगती है

  • “raw assembly vs cooked assembly” वाली अभिव्यक्ति दिलचस्प लगी।
    पहले CPU के लिए assembly code हाथ से लिखना आम बात थी। GPU के मामले में भी इससे ज़्यादा डरने की ज़रूरत नहीं है। आखिरकार compiler भी वही code बनाता है

  • गणितीय रूप से देखें तो inference मूलतः बुनियादी linear algebra/BLAS operations ही है।
    सोचता हूँ क्या dtype और sparsity को ध्यान में रखकर 80% use cases को cover करने वाला कोई सरल inference API बनाया जा सकता है। शायद उसे CUDA जितना जटिल होने की ज़रूरत नहीं होगी

  • composable-kernel मेरे Gentoo system पर सबसे ज़्यादा OOM (memory shortage) कराने वाला software रहा है। Clang जब CK compile करता है, तो प्रति thread 2.5GB से ज़्यादा इस्तेमाल करता है

    • मैंने Debian के लिए CK package की समीक्षा की थी, और -j32 पर build करते समय 64GB workstation पर भी OOM हो गया। -j1 पर चलाया तो 190 घंटे में जाकर सफल हुआ।
      official build servers पर शायद architectures की संख्या घटानी पड़ेगी। सुना है ज़्यादातर dependency packages में सिर्फ headers की ज़रूरत होती है। build time कम करने के लिए सुधार का काम चल रहा है
    • एक single kernel के लिए template instantiation में 10 मिनट से ज़्यादा लगना हैरान करने वाला स्तर है।
      device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, समझ नहीं आता clang से आखिर कराया क्या जा रहा है

  • सोच रहा हूँ क्या इस तरह की प्रगति consumer AMD chips पर भी LLM को अच्छी तरह चलाने में मदद करेगी।
    उदाहरण के लिए, मैं AMD CPU/iGPU और RTX 5080 वाले HP OMEN MAX 16 laptop पर विचार कर रहा हूँ। क्या AMD पक्ष RTX से प्रतिस्पर्धा कर पाएगा?

    • हम मान लेते हैं कि dGPU हमेशा तेज़ होगा, लेकिन memory capacity limits बड़ी रुकावट बनती हैं।
      यह blog post या high-end Mac के नतीजे देखना दिलचस्प है
    • मैं RTX7900XTX पर Ollama के साथ Qwen3 Coder 30B चला रहा हूँ। यह काफ़ी अच्छा काम करता है। लगता है कुछ load system memory और Ryzen 7 CPU पर चला जाता है।
      यह Sonnet 4 API से थोड़ा धीमा है, लेकिन local में इस स्तर की performance मेरे लिए काफ़ी संतोषजनक है।
      Opencode + Ollama + Qwen3 Coder का संयोजन ClaudeCode + Sonnet4 का शानदार विकल्प है।
      हाँ, अगर AI से पूरा coding काम कराना हो तो राय अलग हो सकती है। लेकिन personal assistant के रूप में यह बेहतरीन है

  • समझ नहीं आता AMD ने B300 को पूरी तरह नज़रअंदाज़ क्यों किया