GPU के प्रदर्शन में विस्फोटक सुधार

 (hazyresearch.stanford.edu)
1 पॉइंट द्वारा GN⁺ 2024-05-13 | 1 टिप्पणियां | WhatsApp पर शेयर करें
  • AI computation की लागत बढ़ने के बीच Hazy Research ने निष्कर्ष निकाला है कि GPU performance optimization की कुंजी NVIDIA H100 के tensor core को लगातार व्यस्त रखना है
  • H100 half-precision matrix multiplication में 989 TFLOPs देता है, लेकिन सामान्य operations में केवल लगभग 60 TFLOPs तक रह जाता है, इसलिए जैसे ही tensor core रुकता है, utilization काफी गिर जाता है
  • Peak performance के करीब पहुंचने के लिए WGMMA, shared memory layout, address generation और occupancy को साथ में संभालना पड़ता है; wgmma.mma_async के बिना microbenchmark में performance peak के लगभग 63% पर ठहर जाती है
  • सार्वजनिक CUDA-embedded DSL ThunderKittens tile और vector abstractions के जरिए swizzling और register layout जैसी जटिलताओं को wrap करता है, जिससे FlashAttention family के kernels लिखना सरल हो जाता है
  • H100 के लिए FlashAttention-2 forward kernel लगभग 100 lines में लिखा जाता है और FlashAttention-2 से लगभग 30% तेज है; Based linear attention kernel 215 TFLOPs पर चलता है

H100 के प्रदर्शन को तय करने वाली शर्तें

  • AI बहुत अधिक compute का उपयोग करता है, और Hazy Research पिछले कुछ वर्षों से AI को कम compute में चलाने या दिए गए compute पर अधिक efficiently चलाने के काम में लगा रहा है
  • व्यावहारिक लक्ष्य GPU को तेज बनाने से सीखी गई बातों को संक्षेप में रखना और तेज kernels लिखने में मदद करने वाला CUDA-embedded DSL ThunderKittens जारी करना है
  • व्यापक रूप से, यह भी बताया गया है कि hardware को समझने से AI compute को देखने का तरीका कैसे बदला

NVIDIA H100 की संरचना और bottlenecks

  • H100 SXM GPU पर चर्चा निम्न configuration के आधार पर की गई है
    • 80GB HBM3, bandwidth 3TB/s
    • 50MB L2 cache, bandwidth 12TB/s, पूरे GPU में 25MB के 2 sections में बंटा और crossbar से जुड़ा

  • 132 SM

    • हर SM में अधिकतम 227KB shared memory सहित 256KB L1 cache होता है, और साथ में लगभग 33TB/s bandwidth मिलती है
    • Hopper का नया hardware Tensor Memory Accelerator(TMA) asynchronous address generation और memory fetch संभालता है
    • हर SM 4 quadrants से बना होता है, और हर quadrant में warp scheduler, 512 vector registers, matrix multiplication के लिए tensor core, और parallel built-in instructions होते हैं
    • सारा compute SM में होता है, और अधिकतर register में process होता है
    • H100 पर performance निकालने की कुंजी tensor core को लगातार fed state में रखना है
    • H100 half-precision matrix multiplication में 989 TFLOPs और “बाकी” operations में लगभग 60 TFLOPs देता है
    • जिन cycles में tensor core इस्तेमाल होता है, उनमें hardware utilization कम से कम 94% तक पहुंचता है
    • जिन cycles में tensor core इस्तेमाल नहीं होता, उनमें utilization अधिकतम 6% तक रहता है

WGMMA: जरूरी लेकिन मुश्किल instruction

  • H100 में warp group matrix multiply accumulate instruction wgmma.mma_async है
    • PTX में wgmma.mma_async
    • SASS में HGMMA / IGMMA / QGMMA / BGMMA
  • पिछले GPUs के wmma.mma.sync, mma.sync synchronous तरीके थे, जिनमें 32 threads वाला एक warp tensor core में data डालता था और result का इंतजार करता था
  • wgmma.mma_async लगातार 128 threads को SM के सभी quadrants में cooperatively synchronize कराता है और shared memory से सीधे asynchronous matrix multiplication शुरू करता है
    • Matrix multiplication चलते समय warps register में दूसरा काम कर सकते हैं
    • Result का इंतजार मनचाहे समय पर किया जा सकता है
  • Microbenchmark में H100 का पूरा compute निकालने के लिए इन instructions की जरूरत पड़ी
    • इनके बिना GPU peak utilization के लगभग 63% पर ठहरता हुआ देखा गया
    • इसकी वजह यह हो सकती है कि tensor core को local resources के लिए भी गहरी hardware pipeline चाहिए
  • सबसे बड़ी कठिनाई memory layout की जटिलता है
    • Unswizzled shared memory layout में coalescing बहुत खराब होती है, जिससे L2 bandwidth की बहुत मांग होती है
    • Swizzled layout की documentation गलत थी, इसलिए इसे समझने में समय लगा
    • Swizzled layout कुछ खास matrix shapes में ही काम करता दिखता है और wgmma.mma_async की अन्य capabilities के साथ अच्छी तरह मेल नहीं खाता
    • Hardware tensor core तक जाते समय sub-matrix transpose कर सकता है, लेकिन केवल तब जब layout swizzled न हो
  • FlashAttention जैसे kernels में TMA और L2 cache पर्याप्त तेज हैं, इसलिए इस समस्या को कुछ हद तक छिपाया जा सकता है
  • Hardware का पूरा उपयोग करने के लिए memory requests को coalesce करना और bank conflict से बचना जरूरी है, इसलिए layout control अहम है

Shared memory और bank conflict

  • Shared memory की single-access latency लगभग 30 cycles लगती है, और इस दौरान SM का tensor core लगभग दो 32x32 square matrix multiplications कर सकता है
  • FlashAttention जैसे पुराने कामों में मुख्य रूप से HBM-SRAM bottleneck पर ध्यान था, और पहले यह bottleneck सच में महत्वपूर्ण था
  • HBM के तेज होने और tensor core के chip के बाकी हिस्सों की तुलना में और तेज बढ़ने के साथ, shared memory की छोटी latency को भी हटाना या छिपाना पड़ता है
  • Shared memory 32 banks में बंटी होती है, इसलिए सावधानी न रखने पर bank conflict होता है
    • अगर एक ही memory bank में कई अलग-अलग memory pieces को साथ में request किया जाए, तो requests serialized हो जाती हैं
    • अनुभव के आधार पर kernel असंतुलित तरीके से धीमा हो सकता है
    • WGMMA और MMA instructions जिन register layouts की मांग करते हैं, उन्हें सरल तरीके से इस्तेमाल करने पर bank conflict हो सकता है
  • समाधान यह है कि अलग-अलग swizzling patterns से shared memory को rearrange करके conflicts से बचा जाए
  • जहां संभव हो, register और shared memory के बीच movement से बचना, और जरूरत पड़ने पर WGMMA और TMA जैसे built-in hardware से asynchronous data movement करना बेहतर है
  • वास्तविक warp का उपयोग करके synchronous movement सबसे सामान्य है, लेकिन सबसे खराब fallback के करीब है

Address generation और TMA

  • H100 में tensor core और memory दोनों तेज हैं, इसलिए fetch किए जाने वाले memory address को generate करने का काम ही chip resources का बड़ा हिस्सा लेता है
    • Complex interleaved patterns या swizzling patterns जुड़ने पर यह और स्पष्ट होता है
  • NVIDIA का Tensor Memory Accelerator(TMA) global/shared memory के multi-dimensional tensor layout को specify करने, उस tensor के subtile को asynchronously fetch करने और completion पर barrier trigger करने की सुविधा देता है
  • TMA address generation cost घटाता है और pipeline बनाना भी आसान करता है
  • TMA को wgmma.mma_async की तरह H100 की potential निकालने के लिए जरूरी माना गया है
    • अनुभव के आधार पर यह WGMMA से भी ज्यादा महत्वपूर्ण हो सकता है
    • यह register resources और instruction dispatch बचाता है
    • इसमें global memory पर asynchronous reduction करने की capability भी है, जो जटिल backward kernels में उपयोगी है
  • TMA में भी swizzling mode समझने के लिए कुछ reverse engineering की जरूरत पड़ी, लेकिन यह WGMMA जितना दर्दनाक नहीं था

Occupancy जो cost छिपाती है

  • CUDA में occupancy का मतलब समान execution hardware पर co-scheduled threads की संख्या है
  • SM quadrant का warp scheduler हर cycle में उस warp को instruction issue करने की कोशिश करता है जो instruction लेने के लिए ready हो
  • H100 पिछली generations की तुलना में occupancy पर कुछ कम निर्भर है
    • Asynchronous capabilities के कारण single instruction stream भी memory fetch, matrix multiply, shared memory reduction और register math को साथ-साथ busy रख सकता है
  • लेकिन occupancy गलतियों और synchronization cost को छिपाने में बहुत उपयोगी है
    • पूरी तरह designed pipeline अतिरिक्त occupancy के बिना भी तेज हो सकती है
    • वास्तविक observations में NVIDIA GPU occupancy को ध्यान में रखकर design किए गए दिखे
    • Synchronization और गलतियों की संभावना ज्यादा होने के कारण occupancy बढ़ाने से realized hardware utilization अक्सर बेहतर हुई
  • H100 में occupancy उपयोगी स्तर पर है, लेकिन A100 और RTX 4090 में यह क्रमशः और महत्वपूर्ण मानी गई
    • संभवतः इसकी वजह H100 की तुलना में synchronous instruction dispatch पर अधिक निर्भरता है

ThunderKittens: CUDA के अंदर छोटा DSL

  • ThunderKittens H100 पर fast kernels आसानी से लिखने के लिए बनाया गया CUDA-embedded DSL है
  • शुरुआत में इसे lab internal use के लिए बनाया गया था, बाद में public किया गया
  • नाम इसलिए रखा गया क्योंकि kittens प्यारे होते हैं और code में kittens:: type करना मजेदार लगा
  • ThunderKittens simplicity को लक्ष्य बनाता है और चार templated types देता है
    • Register tiles: register file पर 2D tensor
    • Register vectors: register file पर 1D tensor
    • Shared tiles: shared memory में 2D tensor
    • Shared vectors: shared memory में 1D tensor
  • Tile को height, width और layout से parameterize किया जाता है
  • Register vector को length और layout से parameterize किया जाता है, और shared vector केवल length का उपयोग करता है
    • Shared vector में आमतौर पर bank conflict नहीं होता
  • दिए गए operations warp level या cooperative warp group level पर tile/vector को manipulate करते हैं
    • initializer: shared vector को zero बनाना आदि
    • unary op: exp आदि
    • binary op: mul आदि
    • row/column op: row_sum आदि
  • ThunderKittens CUDA के अंदर embedded है, इसलिए Triton जैसी libraries के उलट abstraction “gracefully” fail करता है
    • कोई feature missing हो तो उसे मनचाहे तरीके से extend किया जा सकता है

FlashAttention के उदाहरण और performance

  • ThunderKittens के उदाहरण के रूप में RTX 4090 के लिए एक simple forward FlashAttention kernel दिया गया
    • केवल headdim=64 handle करता है
    • n 256 का multiple होना चाहिए
    • लगभग 60 lines CUDA code में लिखा गया
    • Hardware utilization 75% है
    • Complexity का अधिकतर हिस्सा swizzling pattern या register layout में नहीं, बल्कि algorithm में है
  • H100 के लिए FlashAttention-2 forward pass भी ThunderKittens से लिखा गया
    • TMA, WGMMA, swizzling mode और descriptor की complexity को ThunderKittens wrap करता है
    • Kernel लगभग 100 lines का है
    • H100 पर FlashAttention-2 से लगभग 30% तेज है
  • ThunderKittens GPU पर इस्तेमाल हो सकने वाले “mini-pytorch” की तरह layouts और instructions को wrap करता है और primitives देता है
  • Based linear attention और आगे जारी होने वाले अन्य architectures के kernels भी साथ में public किए गए
    • Based linear attention kernel 215 TFLOPs पर चलता है
    • Algorithm के अपने recompute को ध्यान में रखें तो यह 300 TFLOPs से ऊपर जाता है
    • Linear attention theoretical रूप से ज्यादा efficient है, लेकिन वास्तविक hardware पर ऐतिहासिक रूप से इसकी efficiency काफी कम रही है
    • माना गया है कि यह result high-throughput applications का दायरा बढ़ा सकता है

Tile-केंद्रित सोच

  • ThunderKittens अच्छी तरह इसलिए काम करता है क्योंकि यह सब कुछ करने की कोशिश नहीं करता
    • CUDA ThunderKittens से कहीं ज्यादा expressive है
    • ThunderKittens छोटा और simple DSL है
  • मुख्य abstraction small tile है, और इसे AI और hardware की दिशा के अनुरूप माना गया है
  • ThunderKittens 16 से छोटी dimensions support नहीं करता
    • माना गया है कि hardware भी ऐसी छोटी dimensions को खास तौर पर नहीं चाहता
    • सवाल इस तरह रखा गया कि “अगर matrix multiply 16x16 से छोटा है, तो क्या आप यकीन से कह सकते हैं कि वह AI है?”
  • CPU era में 32-bit word को register मानने वाला नजरिया AI hardware के लिए फिट नहीं बैठता
    • CUDA का 1024-bit vector register सही दिशा में एक कदम माना गया है
    • यहां register 16x16 tile का data है
  • AI अब भी matrix multiply, reduction और reshape पर केंद्रित है, इसलिए tile abstraction AI और hardware दोनों के लिए उपयुक्त माना गया है
  • आगे AI ideas को hardware पर अच्छी तरह map होने वाले तरीके से rearrange करना होगा
    • Recurrent state size इतना बड़ा होना चाहिए कि SM में fit हो सके
    • Compute density hardware की जरूरत से कम नहीं होनी चाहिए
    • Hardware से सीखी गई बातों को AI design में align करना आगे की महत्वपूर्ण दिशा है

AMD support plan

  • ThunderKittens में AMD hardware support जल्द आने वाला है

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

1 टिप्पणियां

 
GN⁺ 2024-05-13
Hacker News की राय

  • "अगर matrix multiplication 16x16 से छोटा है, तो क्या वह सच में AI है?" यह सवाल दिलचस्प है
    AI hardware की requirements अब और साफ़ होती जा रही हैं। GPU मूल रूप से बिल्कुल अलग काम के लिए डिज़ाइन किए गए थे, लेकिन matrix multiplication hardware अच्छा होने की वजह से AI में इस्तेमाल हुए, और "AI GPU" असल GPU में मौजूद कुछ features को हटा सकता है
    number representation भी 16-bit floating point, 8-bit, 2-bit, 1-bit की ओर छोटा होता जा रहा है, और किसी बिंदु पर उपयुक्त संतुलन तय हो जाएगा। यह लेख दिखाता है कि 16x16 tile को पसंद करने वाला hardware काफ़ी उचित है। अभी कोई न कोई शायद पहले से ही VHDL में ऐसी चीज़ लिख रहा है, या जल्द ही लिखेगा
    आख़िरकार ऐसा लगता है कि एक अधिक सरल, कम general-purpose और सस्ता device आएगा, जो सिर्फ़ "AI" workloads को बिना ज़्यादा बेकार hardware बोझ के अधिकतम कुशलता से चलाएगा

    • GPU पहले से ही जितना संभव हो उतनी lean AI machine में विकसित होते आए हैं। कम से कम 2014 में Nervana की स्थापना के समय से यह कहा जाता रहा कि GPU पुरानी तकनीक हैं और AI के लिए उपयुक्त नहीं, लेकिन शायद किसी ने नहीं सोचा था कि GPU इतनी तेज़ी से AI machine में बदल जाएँगे
    • Apple कई सालों से पहले ही इस दिशा में जा रहा है। die पर मौजूद NPU, GPU या CPU से पूरी तरह अलग है[1]
      Nvidia शायद इस पर काम कर रहा होगा, लेकिन game/entertainment/crypto/AI को एक साथ बाँधने वाला device, यानी video card वाला form factor बनाए रखना, व्यवसायिक रूप से बेहतर विकल्प हो सकता है
      [1] https://github.com/hollance/neural-engine/blob/master/docs/a...
    • "NVIDIA के झूठ" वाला हिस्सा प्रतिस्पर्धा की गहराई दिखाता है। documentation error पूरी तरह संयोग होना मुश्किल है, और diagrams चुराना या कॉपी करना आसान होता है, इसलिए संभव है कि Nvidia ने जानबूझकर उन्हें वहीं छोड़ा हो
      इससे Naveen Rao की Nervana का वह समय याद आता है जब वह Nvidia के अपने driver से तेज़ Nvidia Maxwell driver बना रही थी। तेज़ी से बढ़ते product में documentation की हर गलती प्रतिस्पर्धात्मक रणनीति नहीं होती, लेकिन यह देखते हुए कि researchers को wgmma का reverse engineering करने में काफ़ी समय लगा, और H100 के आसपास अमेरिका-चीन राजनीतिक स्थिति भी है, ऐसा लगता है कि Nvidia अपनी moat बचाने के लिए पुराने तरीके अपना रहा है
      इसलिए H100 की विशेषताओं में ज़रूरत से ज़्यादा गहराई में जाने के बजाय, यह भी देखना चाहिए कि "AI किस तरह का hardware चाहता है" इस सवाल में व्यावसायिक परिस्थितियाँ भी शामिल हैं
    • AMD पहले ही Versal लाइन की दूसरी पीढ़ी तक पहुँच चुका है
      https://www.amd.com/en/products/accelerators/alveo/v80.html
      XDNA Architecture
      https://www.amd.com/en/technologies/xdna.html
    • क्या Google भी लगभग 10 साल से ऐसे devices नहीं बना रहा है?

  • "NVIDIA के झूठ। असली 128b swizzled wgmma layout के बारे में यह बेहद भ्रामक wording है। इस diagram की वजह से मेरी ज़िंदगी के 3 हफ़्ते ऐसे निकल गए जिन्हें वापस नहीं लाया जा सकता, इसलिए यह सार्वजनिक शर्मिंदा करना बनता है" — यह हिस्सा प्रभावशाली लगा
    AI की प्रगति का बहुत बड़ा हिस्सा matrix multiplication optimization जैसी engineering में है, और उस engineering का काफ़ी हिस्सा NVIDIA chips की reverse engineering है — यह बात कितनों को चौंकाएगी, यह जानने की उत्सुकता है

    • architecture अपने आप में बहुत बड़ा फ़र्क नहीं बनाती। अगर काफ़ी बड़े data पर काफ़ी बड़े model को train किया जाए, तो architecture चाहे जो हो, नतीजे अक्सर मिलते-जुलते आते हैं। इसलिए कहा जा सकता है कि AI की ज़्यादातर प्रगति अब इस वजह से है कि हम matrices को बेहद तेज़ी से multiply कर पा रहे हैं

  • warp scheduler, 4 quadrants, tensor memory accelerator, unswizzled wgmma layout…
    GPU terminology और Star Trek-शैली की technobabble के बीच की सीमा अब और धुंधली होती जा रही है

    • लेख पढ़ते समय मुझे इसका कुछ अंदाज़ा था, लेकिन "quadrant को tensor accelerator में warp करना" जैसा वाक्य सचमुच Star Trek जैसा लगता है
      दूसरे लेख पढ़ते हुए भी कभी-कभी ऐसा ख़याल आया है। अगर किसी को यहाँ के लेखों का लिंक देकर पढ़ने को कहा जाए, तो उसे कैसा लगेगा? शायद जैसे वह warp core पर चर्चा कर रहे Trek fans के किसी सम्मेलन में आ गया हो
    • यह पढ़कर मैं एक कदम पीछे हटकर इन शब्दों को नई नज़र से देखने लगा, और बात इतनी सही लगी कि हँसी आ गई

  • AI inference की बिजली खपत घटाने और speed बढ़ाने के लिए analog approximation circuits की ओर जाना सबसे अच्छा हो सकता है
    ज़रूरी यह नहीं कि पूरी तरह सटीक floating-point multiplication और addition हो, बस ऐसा device चाहिए जो दो input voltages लेकर multiplication result के काफ़ी करीब output voltage दे सके

    • मैं एक ऐसे व्यक्ति को जानता हूँ जो इस दिशा में काम करता है; सुना है कि बड़ी चुनौतियाँ हैं मौजूदा chip manufacturing तकनीक से analog logic को संभव बनाने का तरीका, ऐसा design जो antenna की तरह व्यवहार न करे, और यह संभावना कि manufacturing tolerances हर physical chip में अलग हों, इसलिए चलने वाले model को हर chip के लिए finely tune करना पड़ सकता है
      बड़ा फ़ायदा यह है कि float16 को 16 तारों से represent करने के बजाय, एक ही तार के voltage से उस संख्या को represent किया जा सकता है। सैद्धांतिक रूप से float32 से बहुत अधिक precision भी संभव हो सकती है। साथ ही, values को arithmetic logic unit में load किए बिना सीधे connect किया जा सकता है, इसलिए die area और power savings संभावित रूप से कई orders of magnitude तक हो सकते हैं
    • मुझे लगता है कि analog circuits के व्यावहारिक रूप से उपयोगी बनने में अभी समय है, लेकिन जहाँ inaccuracy स्वीकार की जा सकती है वहाँ noisy digital circuits काम आ सकते हैं
      उदाहरण के लिए, अगर दस लाख output bits में से एक bit flip हो जाए और बदले में performance/power ratio बेहतर हो जाए। float32 में यह मुश्किल होगा, क्योंकि एक अकेला infinity value सब बिगाड़ सकता है, लेकिन int8 में जहाँ 0 चाहिए था वहाँ कभी-कभी 128 आ जाए, यह शायद सहन किया जा सकता है
      [1] मुझे नहीं पता कि H100 की matrix floating-point units वास्तव में IEEE 754 का पालन करती हैं या नहीं
    • एक कदम आगे बढ़कर, मुझे लगता है कि कुछ ऐसा चाहिए जो असली biological brain के काम करने के तरीके से मिलता-जुलता हो, लेकिन आसानी से बनाया जा सके
      biological neural networks सामान्य artificial neural networks की तरह लगभग fully connected नहीं होते, और neurons के input/output connection factors 10 से कम होते हैं, इसलिए वे बहुत local होते हैं। जहाँ तक हमें पता है, biology में backpropagation भी नहीं है; उसकी जगह feedback और recurrence है
      यह भी हो सकता है कि कुछ auxiliary cells या processes हों जो central nervous system के उन कार्यों के लिए ज़रूरी हों जिन्हें हम अभी नहीं जानते। ऊँचे स्तर पर भी काफ़ी मात्रा में "hardcoded" connectivity हो सकती है, और इसका कुछ हिस्सा पहले से ज्ञात है। उदाहरण के लिए, कान के auditory neurons जुड़े होते हैं, और sound location पता करने के लिए convolution जैसा कुछ होता है। यह emergent phenomenon नहीं है, बल्कि training के बिना संभव कार्य है
      जीवन ने अरबों वर्षों और लगभग उतनी ही पीढ़ियों में इसे खोज लिया, इसलिए यह आश्चर्यजनक नहीं है। सैद्धांतिक रूप से यह software में भी संभव हो सकता है, लेकिन primate/human brain के एक ट्रिलियन से अधिक neurons को देखते हुए, आज की thousand-core machines पर भी यह बेहद कठिन है। "cloud" भी शायद आवश्यक connectivity और latency को पूरा नहीं कर पाएगा
      अगर इस approach से worm या insect स्तर तक का सफल modeling हो सके, तो वह काफ़ी शानदार होगा
    • पर्याप्त range और precision को एक साथ संतुष्ट करना लगभग असंभव लगता है
    • सच कहूँ तो, debug करने के लिए यह एक बुरा सपना लगता है

  • यह पढ़कर CS 149 parallel programming class में महसूस की गई खुशी फिर याद आ गई

    • Kayvon और Kunle कमाल के हैं। मैंने दो semester पहले CS149 Parallel Programming लिया था और वह सच में बहुत अच्छा था :)

  • इस लेख की writing style सच में प्रभावशाली है, और AMD MI300x पर इसे देखना दिलचस्प होगा। अगर आप चाहते हैं कि मैं अपने hardware पर थोड़ा समय लगाऊँ, तो बताइए

    • जानना चाहता हूँ कि क्या आपने AMD products के साथ AI काम बहुत किया है। मैं RTX 4090 पर 2500 डॉलर से ज़्यादा खर्च नहीं करना चाहता, इसलिए experiment या शुरुआती उपयोग के लिए RX 7900XTX पर विचार कर रहा हूँ
      यह असल में कितना अच्छा काम करेगा, या थोड़ा और बचाकर 7900 XT की जगह XTX लेना बेहतर होगा, और VRAM कम होने से real-world usability पर कितना असर पड़ेगा, यह भी जानना चाहता हूँ
    • अच्छी writing स्पष्ट होनी चाहिए, अस्पष्ट नहीं। बोलते समय आप बीच में रोककर clarification माँग सकते हैं, लेकिन writing को message पहुँचाने का सिर्फ़ एक मौका मिलता है
      पाठक को यह समझने के लिए कि लेखक क्या कहना चाह रहे हैं, knowyourmeme.com तक नहीं जाना चाहिए। मुझे तो यह भी समझ नहीं आता कि इस title का मतलब क्या है, इसलिए मुझे लगता है कि यह लक्ष्य से काफ़ी दूर निकल गया
    • सच में? Wallstreetbets वाले दिनों का PTSD याद आ गया

  • मैं सोच रहा हूँ कि इस तरह की पोस्ट को पूरी तरह समझने के लिए कहाँ से शुरू किया जाए और कौन-सा roadmap follow किया जाए

    • GPU programming सीखने के लिए एक अच्छी lecture series है। लगभग lecture 4.0 तक पहुँचते-पहुँचते ज़रूरी basics मिल जाते हैं: https://youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6Srgd...
      और vector-matrix multiplication करने वाला CUDA kernel खुद लिखकर देखना अच्छा रहेगा। pycuda इस्तेमाल करें तो आप kernel पर ध्यान दे सकते हैं और बाकी Python में लिख सकते हैं। ChatGPT से कहिए कि आप 4000-element vector और 4000x12000 matrix को multiply करने का implementation खुद बनाना चाहते हैं, और वह आपको पूरा process guide करे
      GPU rental के लिए Runpod अच्छा है, और अभी वहाँ low-end GPUs से लेकर H100 तक मिलते हैं। शुरुआत low-tier GPU से की जा सकती है
    • अगर गहराई में जाना है, तो Spiral matrix multiplication playlist देखना अच्छा रहेगा: https://www.youtube.com/playlist?list=PL04PGV4cTuIWT_NXvvZsn...
      मैंने Spiral के साथ matrix multiplication kernels को implement और optimize करने में 2 महीने लगाए थे

  • GitHub README का graph(https://github.com/HazyResearch/ThunderKittens/blob/main/att...) बहुत उलझाने वाला है। क्या ऐसे wavy bar charts सच में जायज़ हैं? :P

    • सहमत। लगता है जैसे किसी optical illusion की कोशिश की गई हो। बेहतर होगा कि bars की बजाय सिर्फ़ numbers ही दिखें
    • यह matplotlib की xkcd theme जैसा लगता है[1]। फिर भी, यह मानता हूँ कि wave effect कुछ ज़्यादा ही है
      [1]: https://matplotlib.org/stable/gallery/showcase/xkcd.html#sph...

  • ThunderKittens नाम कमाल का है। मैं देखना चाहूँगा कि ThunderKittens FlashAttention backward pass को संभालता है, जो forward pass से एक स्तर ज़्यादा कठिन है

  • क्या इस तरह का शोध आज के NPU बनाने वाली टीमें पहले से नहीं कर रही हैं? उदाहरण के लिए, Groq chip AI-specific architecture का उपयोग करता है, इसलिए वह अभी जैसी performance दे पाता है। consumer side में Apple Silicon भी काफ़ी सक्षम है
    मैं इस क्षेत्र का विशेषज्ञ नहीं हूँ, लेकिन ऐसा लगता है कि अपेक्षाकृत धीमे path से communicate करने वाले general-purpose processors के साथ सीमाएँ होंगी। hardware स्तर पर design पर फिर से सोचना, और अंततः consumer market में कीमत कम करना, लंबी अवधि की बेहतर रणनीति लगता है

    • consumer side में Apple Silicon के काफ़ी सक्षम होने वाली बात पर मुझे यक़ीन नहीं है। reddit के localllama subreddit को देखें, वहाँ CPU users की बहुत-सी निराश पोस्टें हैं जो उपयोगी speed पाने की कोशिश कर रहे हैं
      कुछ सौ डॉलर में Nvidia GPU खरीदा जा सकता है, या 900 डॉलर में 4050 6GB VRAM gaming laptop मिल सकता है, इसलिए CPU-based AI को सक्षम कहना मुश्किल है
      मेरे workplace में भी GPU नहीं था, इसलिए CPU-based setup पर कोशिश की, लेकिन छोटे models इस्तेमाल करके इंतज़ार करने के अलावा वह व्यावहारिक नहीं था। अंत में मुझे GPU computer माँगना पड़ा
      "तकनीकी रूप से संभव है" और "वास्तव में उपयोग में अच्छा है" अलग बातें हैं। Nvidia सचमुच इस्तेमाल करने में अच्छा था, और CPU कष्टदायक और निराशाजनक था