
NVIDIA SASS निर्देश शब्दकोश का रिवर्स इंजीनियरिंग, कर्नेल ऑडिट और GPU आर्किटेक्चर में पैटर्न पहचान।
नियंत्रित कर्नेल से लेकर प्रोडक्शन ऑडिट तक NVIDIA SASS का रिवर्स इंजीनियरिंग।
लेख 1 · लेख 2 · ज्ञानकोष · पैटर्न लाइब्रेरी · SM120 निर्देश शब्दकोश · एन्कोडिंग नोट्स · यहाँ से शुरू करें · प्रोजेक्ट संरचना · टेंसर-कोर अध्याय · योगदान
SASS किंग NVIDIA SASS का एक व्यवस्थित रिवर्स-इंजीनियरिंग प्रोजेक्ट है, जो संकलित CUDA बाइनरी के अंदर उत्सर्जित मूल GPU निर्देश सेट है। यह प्रोजेक्ट SM120 / SM120a उपभोक्ता Blackwell हार्डवेयर से शुरू होता है और समय के साथ एक पूर्ण क्रॉस-आर्किटेक्चर ISA और पैटर्न लाइब्रेरी की ओर बढ़ता है।
लक्ष्य व्यावहारिक है: एक कर्नेल इंजीनियर को SASS डंप खोलने, कंपाइलर पैटर्न पहचानने, प्रदर्शन-प्रासंगिक संरचनाओं की पहचान करने, और बाइनरी को स्रोत-स्तरीय अनुकूलन निर्णयों से जोड़ने में मदद करना।
प्रोजेक्ट ने अपनी प्रारंभिक चरण 3 पैटर्न लाइब्रेरी पूरी कर ली है: 29 पुन: प्रयोज्य SASS हस्ताक्षर अब patterns/ के अंतर्गत औपचारिक रूप में हैं, knowledge/FINDINGS.md को पूर्ण साक्ष्य श्रृंखला के रूप में रखा गया है। अगला बड़ा कदम चरण 4 है: वास्तविक प्रोडक्शन कर्नेल पर उन पैटर्न को लागू करना।
रिपॉजिटरी एक साक्ष्य पाइपलाइन के रूप में व्यवस्थित है:
corpus/ नियंत्रित कर्नेल और रॉ SASS साक्ष्य
knowledge/ प्रोजेक्ट-व्यापी निष्कर्ष, निर्देश नोट्स और एन्कोडिंग नोट्स
patterns/ पुन: प्रयोज्य चरण 3 ऑडिट हस्ताक्षर
production/ चरण 4 वास्तविक-कर्नेल ऑडिट
भावना में समान अंतिम व्यापक सार्वजनिक SASS रिवर्स-इंजीनियरिंग कार्य 2018 में वोल्टा और ट्यूरिंग पर जिया एट अल. का था। एम्पीयर, हॉपर और ब्लैकवेल ने निर्देश मिश्रण को काफी हद तक बदल दिया है: एसिंक कॉपी पथ, टेंसर-कोर परिवार, मैट्रिक्स लोड/स्टोर निर्देश, स्पार्स और स्केल्ड MMA रूप, और नए यूनिफ़ॉर्म-रजिस्टर प्रवाह।
SASS किंग नियंत्रित माइक्रो-कर्नेल, रॉ SASS पढ़ना, रनटाइम प्रोब और प्रोडक्शन-कर्नेल ऑडिट को मिलाकर उस अंतर को भरता है।
औपचारिक पैटर्न लाइब्रेरी चरण 3 का मुख्य आउटपुट है। यह अध्याय-स्थानीय साक्ष्य को पुन: प्रयोज्य ऑडिट हस्ताक्षर में बदलता है, ताकि एक ऑडिट हर बार पूर्ण शोध श्रृंखला को फिर से लिखने के बजाय एक नामित पैटर्न का हवाला दे सके।
चरण 3 को पूर्ण माना जाता है क्योंकि:
knowledge/FINDINGS.md में स्रोत साक्ष्य से बंधे रहते हैं;patterns/README.md से शुरू होता है;प्रत्येक पैटर्न पृष्ठ में शामिल हैं:
ऑडिट-सामने सूचकांक के रूप में patterns/README.md का उपयोग करें। जब आपको किसी पैटर्न के पीछे लंबे शोध संदर्भ की आवश्यकता हो, तो knowledge/FINDINGS.md का उपयोग करें।
चरण 3 यह दावा नहीं करता कि प्रत्येक NVIDIA SASS व्यवहार डीकोड किया गया है। यह एक पुन: प्रयोज्य SM120 / SM120a पैटर्न परत स्थापित करता है जो मैन्युअल प्रोडक्शन ऑडिट शुरू करने के लिए पर्याप्त है। रनटाइम लेआउट डीकोड, पूर्ण नियंत्रण-कोड बिट प्लेसमेंट, स्वचालित क्यूबिन रिपोर्टिंग और क्रॉस-आर्किटेक्चर रीप्ले भविष्य के कार्य बने हुए हैं।
सार्वजनिक लेख:
नियंत्रित भिन्नता. दो कर्नेल ठीक एक चर में भिन्न होते हैं: डेटाटाइप, ऑपरेंड ऑर्डर, अनरोल फैक्टर, मेमोरी लेआउट, या संकलन लक्ष्य। SASS अंतर कंपाइलर निर्णय को अलग करता है।
सख्त दावा टैग. प्रत्येक तकनीकी दावा एक टैग का उपयोग करता है:
ऊपर-से-नीचे और नीचे-से-ऊपर एक साथ। माइक्रो-कर्नेल व्यक्तिगत निर्देशों और कंपाइलर निर्णयों को अलग करते हैं। प्रोडक्शन-जैसे कर्नेल दिखाते हैं कि वास्तविक कोड में कौन से पैटर्न मायने रखते हैं।
पैटर्न-प्रथम ऑडिट। एक प्रोडक्शन ऑडिट को दृश्य SASS हस्ताक्षर से मिलान करने और अपने आत्मविश्वास सीमाओं, एंटी-पैटर्न और खुले अंतर को स्थानांतरित करने के बाद ही एक औपचारिक PATTERN-NN पृष्ठ का हवाला देना चाहिए।
पहला पास SM120 टेंसर-कोर और मेमोरी पाइपलाइन पर केंद्रित है:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , प्रोजेक्ट यह दिखावा नहीं करता कि ISA अभी भी पूर्ण है। सार्वजनिक शब्दकोश ट्रैक करता है कि क्या देखा और समझाया गया है; knowledge/encoding/ के तहत गहरे पृष्ठ उन परिवारों को ट्रैक करते हैं जिनमें मैचर-शैली दस्तावेज़ीकरण के लिए पर्याप्त साक्ष्य हैं।
SASS किंग बिट-स्तरीय SASS डिसअसेंबलरों से प्रतिस्पर्धा नहीं करता है। प्रोजेक्ट स्थानीय डंप को प्राथमिक साक्ष्य के रूप में उपयोग करता है और निर्देश क्षेत्रों, शेड्यूलिंग तालिकाओं, प्रेडिकेट्स और रजिस्टर ट्रैकिंग के लिए redplait/denvdis को क्रॉस-चेक के रूप में उपयोग कर सकता है। denvdis निम्न-स्तरीय एन्कोडिंग व्याख्याओं को मान्य कर सकता है; SASS किंग नियंत्रित-भिन्नता साक्ष्य, सिमैंटिक पैटर्न परत और प्रोडक्शन-ऑडिट व्याख्या का मालिक है।
flowchart LR
P1["चरण 1<br/>शिक्षण कर्नेल<br/>01-12"] --> P2["चरण 2<br/>SM120 टेंसर-कोर कोर्पस<br/>13-25"]
P2 --> P25["चरण 2.5<br/>denvdis क्रॉस-सत्यापन<br/>बिट-स्तरीय बैकएंड"]
P25 --> P3["चरण 3<br/>पैटर्न लाइब्रेरी<br/>कंपाइलर हस्ताक्षर"]
P3 --> P4["चरण 4<br/>प्रोडक्शन ऑडिट<br/>वास्तविक कर्नेल"]
P4 --> P5["चरण 5<br/>ऑडिट उपकरण<br/>cubin रिपोर्ट"]
P5 --> P6["चरण 6<br/>क्रॉस-आर्किटेक्चर रीप्ले<br/>SM80/86/89/90a/100a/120"]
classDef done fill:#0b6d55,color:#fff,stroke:#0b6d55;
classDef active fill:#f4c95d,color:#111,stroke:#b89422;
classDef planned fill:#1f2937,color:#fff,stroke:#6b7280;
class P1,P2,P25,P3 done;
class P4 active;
class P5,P6 planned;
कर्नेल 01-12 बेसलाइन SASS अवधारणाओं को स्थापित करते हैं: FMA फ्यूजन, स्कोरबोर्ड व्यवहार, लूप लोअरिंग, शेयर्ड मेमोरी, ग्लोबल मेमोरी, वार्प प्राइमिटिव, स्लो-पथ गणित और स्थानीय-मेमोरी स्पिल।
कर्नेल 13-25 वर्तमान SM120 टेंसर-कोर पथ को कवर करते हैं:
प्रोडक्शन ऑडिट पैटर्न लाइब्रेरी पर निर्भर होने से पहले redplait/denvdis को SM120 / SM120a के लिए बिट-स्तरीय क्रॉस-चेक बैकएंड के रूप में मान्य करें। पास प्रतिनिधि स्थानीय क्यूबिन या डंप पर nvd -O, nvd -S, nvd -p, और जहाँ उपयोगी हो nvd -T चलाता है जो HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR, और विचलन मार्कर को कवर करता है।
आउटपुट knowledge/DENVDIS_INTEGRATION.md है: परिवार से denvdis पहचान स्थिति, संशोधक कवरेज, उजागर नियंत्रण-कोड क्षेत्रों और SASS किंग कार्रवाई के लिए एक तथ्यात्मक संगतता तालिका। denvdis आउटपुट सहायक साक्ष्य है, स्थानीय डंप अवलोकनों का प्रतिस्थापन नहीं।
आवर्ती संरचनाओं को पुन: प्रयोज्य हस्ताक्षर में औपचारिक रूप दिया गया:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMASTSM -> BAR -> LDS -> STGप्रारंभिक चरण 3 लाइब्रेरी में patterns/ के अंतर्गत 29 पैटर्न पृष्ठ हैं। knowledge/FINDINGS.md शोध लॉग और सत्य स्रोत बना हुआ है; patterns/ ऑडिट-सामने प्रवेश बिंदु है।
चरण 3 प्रारंभिक लाइब्रेरी स्तर पर पूर्ण है। रनटाइम लेआउट डीकोड, पूर्ण नियंत्रण-कोड बिट प्लेसमेंट और क्रॉस-आर्किटेक्चर रीप्ले जैसे शेष आइटम अंतर या भविष्य के चरणों के रूप में ट्रैक किए जाते हैं, चरण 4 प्रोडक्शन ऑडिट शुरू करने में अवरोधक नहीं।
FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad और संबंधित परियोजनाओं जैसी लाइब्रेरीज़ के वास्तविक कर्नेल पर पैटर्न लाइब्रेरी लागू करें। लक्ष्य एल्गोरिथमिक पैटर्न द्वारा प्रतिनिधित्वपूर्ण कवरेज है, प्रति कर्नेल एक मार्कडाउन फ़ाइल नहीं।
पहला चरण 4 प्रदेय एक मैन्युअल ऑडिट रिपोर्ट होना चाहिए जो:
PATTERN-NN पृष्ठों का हवाला देता है;एक पाइपलाइन बनाएं जो एक क्यूबिन लेता है, ज्ञात पैटर्न का पता लगाता है, और अनुकूलन-उन्मुख रिपोर्ट जारी करता है।
अतिरिक्त लक्ष्यों पर पद्धति को दोहराएं:
.
├── corpus/ # नियंत्रित कर्नेल, डंप और अध्याय लेख
│ ├── basics/ # कर्नेल 01-08: स्केलर/वेक्टर और मेमोरी मूल बातें
│ ├── warp_collectives/ # कर्नेल 09-10: शफल, वोट, कमी
│ ├── math_and_spills/ # कर्नेल 11-12: स्लो पथ और स्पिल
│ └── tensor_cores/ # कर्नेल 13-25: टेंसर-कोर अध्ययन
├── knowledge/ # निष्कर्ष, शब्दकोश, एन्कोडिंग नोट्स
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # औपचारिक चरण 3 पैटर्न लाइब्रेरी
├── production/ # चरण 4 प्रोडक्शन-कर्नेल ऑडिट
├── docs/ # ऑनबोर्डिंग, संरचना और रिलीज़-सामने नोट्स
└── guide/ # बाहरी SASS पढ़ने की गाइड सबमॉड्यूल
प्रत्येक अध्याय फ़ोल्डर में स्रोत कर्नेल, प्रासंगिक संकलित आर्टिफैक्ट, SASS डंप (जब वे मान्य साक्ष्य सेट का हिस्सा हों) और एक conclusion<N>.md लेख शामिल है।
प्रत्येक निर्देशिका में क्या शामिल है, इसकी पूर्ण व्याख्या के लिए प्रोजेक्ट संरचना पढ़ें।
cuobjdump --dump-sass रॉ डिसअसेंबली के लिए।gpuasm.com स्कोरबोर्ड, स्टॉल, दबाव और निर्भरता तीरों के लिए।%clock माइक्रोबेंचमार्क निर्देश विलंबता जांच के लिए।nvcc -Xptxas -v रजिस्टर और स्पिल मेटाडेटा के लिए।SASS किंग एल्गोरिथमिक पैटर्न परत पर काम करता है: संकलित कर्नेल कैसे संरचित होते हैं, इसे पहचानना और उन संरचनाओं को स्रोत-स्तरीय अनुकूलन निर्णयों से जोड़ना।
योगदान का स्वागत है, विशेष रूप से:
अपेक्षित मेटाडेटा और लेखन मानक के लिए CONTRIBUTING.md देखें।
फ्लोरियन मट्टाना। florianmattana.com
| यदि आप चाहते हैं... | यहाँ से शुरू करें | फिर पढ़ें |
|---|
| 10 मिनट में प्रोजेक्ट समझें | docs/README.md | docs/START_HERE.md, फिर docs/PROJECT_STRUCTURE.md |
| साक्ष्य को पुन: उत्पन्न करें | corpus/README.md | एक अध्याय conclusion*.md, फिर इसका .sass डंप |
| सत्य का स्रोत खोजें | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| किसी नए डंप में एक पैटर्न पहचानें | patterns/README.md | संबंधित patterns/NN-*.md पृष्ठ |
| प्रोडक्शन ऑडिट शुरू करें | production/README.md | संबंधित PATTERN-NN पृष्ठ और स्रोत साक्ष्य |
| कोई सुधार या डंप योगदान करें | CONTRIBUTING.md | docs/START_HERE.md |
| क्षेत्र | स्थिति | कहाँ |
|---|
| SM120 शिक्षण कर्नेल | कर्नेल 01-12 के माध्यम से पूर्ण | corpus/basics/01_vector_add/ से corpus/math_and_spills/12_register_spill/ तक |
| टेंसर-कोर अध्ययन | कर्नेल 25 के माध्यम से पूर्ण | corpus/tensor_cores/ |
| वैश्विक निष्कर्ष | सक्रिय सत्य स्रोत | knowledge/FINDINGS.md |
| SM120 निर्देश शब्दकोश | सक्रिय, साक्ष्य-समर्थित | knowledge/SASS_INSTRUCTIONS_SM120.md |
| एन्कोडिंग पायलट | LDSM, STSM, QMMA से शुरू | knowledge/encoding/ |
| denvdis क्रॉस-सत्यापन | प्रारंभिक पास पूर्ण; गहरा नियंत्रण-कोड अंतर शेष | knowledge/DENVDIS_INTEGRATION.md |
| पैटर्न लाइब्रेरी | प्रारंभिक चरण 3 लाइब्रेरी पूर्ण | patterns/ |
| प्रोडक्शन ऑडिट | अगला चरण | production/ |
| पैटर्न परिवार | उदाहरण | कहाँ |
|---|
| टेंसर-कोर गणना | HMMA, QMMA, OMMA संचायक श्रृंखलाएँ; स्पार्स मेटाडेटा; संकीर्ण खंड | patterns/02-* से patterns/04-*, patterns/10-*, patterns/21-* |
| मैट्रिक्स मेमोरी और एपिलॉग | LDSM, STSM, एसिंक कॉपी पाइपलाइन, REDG कमी एपिलॉग | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| नियंत्रण प्रवाह | विचलन/पुनर्मिलन, लूप बैक-एज, प्रेडिकेटेड एग्जिट, कोल्ड ट्रैप, स्थानीय CALLs | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| मेमोरी और रजिस्टर | वेक्टराइज्ड ग्लोबल मेमोरी, स्पिल, शेयर्ड-मेमोरी स्टेजिंग, डिस्क्रिप्टर, यूनिफ़ॉर्म-रजिस्टर प्रवाह | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| अंकगणित और शेड्यूलिंग | FFMA फ्यूजन, स्थिरांक, MUFU स्लो-पथ, स्कोरबोर्ड, लाइफटाइम रिसाइक्लिंग | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| वार्प सामूहिक | वार्प कमी, शफल/वोट/मैच/सिंक प्राइमिटिव | patterns/01-*, patterns/25-* |
| टैग | अर्थ |
|---|
[OBS] | सीधे डंप, लॉग, रनटाइम आउटपुट या प्रोफाइल में देखा गया। |
[INF] | देखे गए साक्ष्य से अनुमानित। |
[HYP] | प्रशंसनीय लेकिन पुष्टि नहीं हुई। |
[RES] | पिछली परिकल्पना बाद के साक्ष्य द्वारा हल की गई। |
[GAP] | स्पष्ट रूप से दस्तावेजीकृत खुला प्रश्न। |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| चरण | स्थिति | आउटपुट | यह क्यों मायने रखता है |
|---|
| 1. शिक्षण कर्नेल | हो गया | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | नियंत्रित CUDA-से-SASS प्रयोगों से पढ़ने की शब्दावली स्थापित करता है। |
| 2. SM120 टेंसर-कोर कोर्पस | हो गया | corpus/tensor_cores/13_hmma_fp16/ से 25_stsm_epilogue/ तक | पहला SM120 / SM120a टेंसर-कोर, मैट्रिक्स-मेमोरी, नियंत्रण-प्रवाह और एपिलॉग साक्ष्य सेट कैप्चर करता है। |
| 2.5. denvdis क्रॉस-सत्यापन | प्रारंभिक पास पूर्ण | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | स्थानीय डंप साक्ष्य को बदले बिना बिट-स्तरीय क्रॉस-चेक के रूप में denvdis का उपयोग करता है। पूर्ण स्टॉल/यील्ड बिट प्लेसमेंट खुला रहता है। |
| 3. पैटर्न लाइब्रेरी | प्रारंभिक लाइब्रेरी पूर्ण | patterns/ | आवर्ती कंपाइलर/SASS संरचनाओं को पुन: प्रयोज्य हस्ताक्षर में बदलता है। |
| 4. प्रोडक्शन ऑडिट | अगला | production/ | परीक्षण करता है कि क्या कोर्पस पैटर्न प्रोडक्शन लाइब्रेरीज़ के वास्तविक कर्नेल की व्याख्या करते हैं। |
| 5. ऑडिट उपकरण | योजनाबद्ध | क्यूबिन-से-रिपोर्ट पाइपलाइन | पैटर्न परत को स्क्रिप्टेबल और दोहराने योग्य बनाता है। |
| 6. क्रॉस-आर्किटेक्चर रीप्ले | योजनाबद्ध | SM80, SM86, SM89, SM90a, SM100a, SM120 तुलना | आर्किटेक्चर-विशिष्ट तथ्यों को सामान्य NVIDIA SASS व्यवहार से अलग करता है। |
| कर्नेल | विषय |
|---|
| 13 | HMMA बेसलाइन, रजिस्टर आवंटन, संचायक श्रृंखलाबद्ध |
| 14 | QMMA FP8 / FP6 / FP4 बेसलाइन |
| 15 | संकीर्ण MMA वेरिएंट |
| 16 | FP4 पीक और ब्लॉक-स्केल्ड OMMA/QMMA |
| 17 | LDSM और मैट्रिक्स-लोड व्यवहार |
| 18 | पाइपलाइन्ड MMA टाइल और एसिंक कॉपी स्टेजिंग |
| 19 | स्पार्स MMA मेटाडेटा |
| 20 | नियंत्रण प्रवाह और बैक-एज डिटेक्शन |
| 21 | विचलन और पुनर्मिलन |
| 22 | STSM मैट्रिक्स-स्टोर व्यवहार |
| 23 | FP4 / FP6 फ़्रैगमेंट लेआउट प्रोब |
| 24 | प्रोडक्शन मिनी-GEMM ऑडिट |
| 25 | STSM एपिलॉग लेआउट और स्टोरबैक सिमैंटिक्स |
| आर्क | प्रतिनिधि GPU | क्यों |
|---|
| SM80 | A100 | डेटासेंटर Ampere बेसलाइन |
| SM86 | RTX 3090 | उपभोक्ता Ampere कोर्पस |
| SM89 | RTX 4090 | सामान्य उपभोक्ता अनुमान कार्ड |
| SM90a | H100 | TMA, WGMMA, वार्प स्पेशलाइज़ेशन, क्लस्टर |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | उपभोक्ता Blackwell प्रारंभिक बिंदु |