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

 (hazyresearch.stanford.edu)
1 पॉइंट द्वारा GN⁺ 2024-05-13 | 1 टिप्पणियां | WhatsApp पर शेयर करें

H100 GPU की विशेषताएँ

  • 80GB HBM3 मेमोरी और 3TB/s बैंडविड्थ प्रदान करता है (वास्तव में यह थोड़ा कम है)
  • 50MB L2 cache और 12TB/s बैंडविड्थ प्रदान करता है। यह GPU में 25MB के दो सेक्शन में विभाजित है और crossbar से जुड़ा है (crossbar प्रदर्शन गिरने का कारण बनता है)
  • 132 Streaming Multiprocessor(SM) से बना है, और प्रत्येक SM में निम्न संरचना होती है:
    • 256KB L1 cache के भीतर अधिकतम 227KB shared memory प्रदान करता है (मिलाकर लगभग 33TB/s बैंडविड्थ)
    • asynchronous address generation और memory fetching सक्षम करने वाला Tensor Memory Accelerator(TMA) प्रदान करता है। on-chip memory network support जैसी सुविधाएँ भी देता है, लेकिन इस पोस्ट में उन्हें शामिल नहीं किया गया है
    • 4 quadrant में विभाजित, और प्रत्येक quadrant में warp scheduler, 512 vector register (प्रत्येक में 32 4-byte word शामिल), matrix multiply के लिए tensor core, और sum/multiply जैसी parallel operations को support करने वाले built-in instruction शामिल हैं

H100 GPU प्रदर्शन अनुकूलन टिप्स

  • Tensor Core का अधिकतम उपयोग करना महत्वपूर्ण है। H100 में 989 TFLOPs का FP16 matrix multiplication प्रदर्शन है, इसलिए Tensor Core utilization के आधार पर पूरे GPU utilization पर बड़ा प्रभाव पड़ता है
  • हालांकि, Tensor Core का अधिकतम उपयोग करने के लिए नीचे की बातों पर ध्यान देना होगा:
    • Warp Group Matrix Multiply Accumulate(WGMMA) instruction का उपयोग आवश्यक है, लेकिन इसका उपयोग कठिन है
    • Shared Memory अपेक्षा से उतनी तेज़ नहीं है, और इसका उपयोग करते समय बहुत सावधानी की ज़रूरत होती है
    • address generation की लागत अधिक है, इसलिए Tensor Memory Accelerator(TMA) आदि का उपयोग कर optimization करना चाहिए
    • Occupancy बढ़ाना अब भी मददगार है, और register मुख्य resource है
  • ये विशेषताएँ केवल H100 पर ही नहीं बल्कि अन्य GPU पर भी कुछ हद तक लागू होती हैं, लेकिन खासकर H100 में Tensor Core का उपयोग महत्वपूर्ण और कठिन दोनों है

ThunderKittens : H100 के लिए अनुकूलित CUDA embedded DSL

  • NVIDIA H100 जैसे नवीनतम GPU का प्रदर्शन अधिकतम करने के लिए विकसित CUDA-आधारित embedded DSL
  • नीचे दिए गए 4 tile-आधारित template type प्रदान करता है:
    • Register Tile (register में संग्रहित 2D tensor)
    • Register Vector (register में संग्रहित 1D tensor)
    • Shared Tile (Shared Memory में संग्रहित 2D tensor)
    • Shared Vector (Shared Memory में संग्रहित 1D tensor)
  • साथ ही tile manipulation के लिए विभिन्न operator (exp, mul, sum आदि) warp या warp group स्तर पर प्रदान करता है
  • मौजूदा Flash Attention, Flash Attention 2 kernel को ThunderKittens में लागू करने पर code काफी सरल हो गया, और H100 पर अधिकतम 30% प्रदर्शन सुधार मिला
  • Based Linear Attention kernel भी ThunderKittens में लागू कर 215 TFLOPs प्रदर्शन हासिल किया गया (algorithm की विशेषता के कारण recompute शामिल करें तो 300 TFLOPs से अधिक)

दार्शनिक दृष्टिकोण से विचार

  • ThunderKittens के अच्छी तरह काम करने का कारण यह है कि यह सब कुछ support करने की कोशिश नहीं करता, बल्कि सरल और GPU architecture के अनुरूप tile-आधारित abstraction प्रदान करता है
  • पारंपरिक 32-bit word के बजाय 1024-bit vector register का उपयोग भी प्रगति है, लेकिन 16x16 tile को register की इकाई के रूप में देखने वाले paradigm shift की ज़रूरत है
  • AI workload अंततः matrix multiplication, reduction, reshape पर केंद्रित होते हैं, इस दृष्टि से tile-आधारित approach उचित है, और hardware दृष्टिकोण से भी systolic array के अलावा छोटे matrix multiplication को support करने की दिशा में विकास होगा
  • आगे बढ़कर AI algorithm को hardware-optimized रूप में design करने की दिशा में सोच बदलने की ज़रूरत है। उदाहरण के लिए, RNN की state size को SM में समाने लायक सीमित करना और computation density को hardware की अपेक्षित स्तर के अनुसार समायोजित करना जैसी approach आवश्यक है

GN⁺ की राय

  • ThunderKittens, CUDA से परिचित developers के लिए एक आकर्षक विकल्प हो सकता है। क्योंकि मौजूदा kernel code को बहुत बदले बिना भी आसानी से प्रदर्शन सुधार हासिल किया जा सकता है
  • हालांकि beginners के लिए अब भी प्रवेश बाधा ऊँची हो सकती है। आगे और अधिक example code और learning material का समर्थन आवश्यक लगता है
  • H100 ही नहीं बल्कि AMD जैसे अन्य GPU तक ThunderKittens support का विस्तार करने की योजना दिलचस्प है। इससे vendor lock-in कम करने में योगदान मिल सकता है
  • अंततः AI model/algorithm को ही hardware-optimized रूप में design करना बहुत महत्वपूर्ण बिंदु है। इसके लिए hardware की विशेषताओं की गहरी समझ पहले से आवश्यक है, और ThunderKittens developers को ऐसी insight पाने में मदद कर सकता है
  • Hazy Research का निरंतर R&D और open source योगदान CUDA ecosystem को सक्रिय बनाने में बहुत मददगार होगा। हालांकि लंबी अवधि में अधिक उच्च abstraction level वाले framework के उभरने की भी आवश्यकता दिखती है

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

1 टिप्पणियां

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

  • AI हार्डवेयर की आवश्यकताएँ धीरे-धीरे अधिक स्पष्ट होती जा रही हैं। GPU मूल रूप से किसी और उद्देश्य के लिए डिज़ाइन किए गए थे, लेकिन उनमें अच्छा matrix multiplication hardware होने के कारण उनका उपयोग AI में हो रहा है। "AI GPU" वास्तविक GPU की कुछ सुविधाओं को हटा सकता है, और रुझान छोटे नंबरों (16-bit, 8-bit, 2-bit, 1-bit floating point) की ओर बढ़ रहा है। यह पेपर संकेत देता है कि 16x16 tile पसंद करने वाले hardware के कई फायदे हैं.

  • AI-समर्पित सह-प्रोसेसर (NPU) की आवश्यकता है। खासकर developers, experts, gamers आदि के लिए prosumer-grade desktop systems में इसकी ज़रूरत है। GPU enterprise में तो काम करते हैं, लेकिन personal computing के संदर्भ में AI के लिए उनका उपयोग करना असुविधाजनक है। खासतौर पर VRAM की सीमा और Vulkan के अलावा standard open API की अनुपस्थिति समस्या है.

  • AI research को आगे बढ़ाने के लिए neuroscience और psychology का बेहतर अध्ययन करना होगा। इसके अलावा, neural networks की graph topology से जुड़ी चीज़ें भी प्रासंगिक हो सकती हैं.

  • क्या यह CUTLASS को user-friendly बनाता है?

  • ThunderKittens का mascot बिल्ली/Sony Aibo जैसा vibe देता है। लगता है कि इसे AI से अच्छी तरह generate किया गया है.

  • अगर Universal Basic Compute (UBC) को Universal Basic Income के विकल्प के रूप में देखा जाए, तो वह बहुत dystopian भविष्य होगा। कल्पना कीजिए कि Nvidia जैसी एक ही कंपनी सारा computing बना रही हो।