Skip to content
KitploitKITPLOIT
उपकरणब्लॉग
जमा करें
उपकरणब्लॉग
जमा करें

हैकिंग, पेनटेस्ट और साइबर सुरक्षा उपकरण आपके सुरक्षा शस्त्रागार के लिए!

Kitploit हैकिंग, साइबर सुरक्षा और पेंटेस्टिंग टूल्स की एक निर्देशिका है। कमजोरियों को खोजने, सिस्टम का विश्लेषण करने, परीक्षण को स्वचालित करने और अपनी सुरक्षा को मजबूत करने के लिए नवीनतम प्रोजेक्ट अपडेट खोजें।

··फ़ीड·संपर्क·गोपनीयता·© 2026 Kitploit

टूल निर्देशिका

श्रेणियाँ

सभी श्रेणियाँ देखें
Loading categories
sass-king — NVIDIA SASS निर्देश शब्दकोश का रिवर्स इंजीनियरिंग, कर्नेल ऑडिट और GPU आर्किटेक्चर में पैटर्न पहचान। | Kitploit
उपकरण/GitHubGitHub/florianmattana/sass-king
एम्बेडेड सिस्टम सुरक्षास्थैतिक विश्लेषणकोड विश्लेषणरिवर्स इंजीनियरिंगहार्डवेयर सुरक्षाहार्डवेयर और IoT सुरक्षाबाइनरी विश्लेषणपेपर और शोधलर्निंग और शिक्षाचयनित संसाधनफर्मवेयर विश्लेषण
317153 महीने पहलेKitploit द्वारा समीक्षित

सबसे लोकप्रिय

सभी देखें →

हमारे समुदाय द्वारा सबसे अधिक उपयोग किए जाने वाले उपकरण खोजें।

सभी उपकरण खोजें

हमारे उपकरणों का संग्रह ब्राउज़ करें

सभी उपकरण देखें →
साझा करें
GitHub
florianmattana/sass-king

sass-king

NVIDIA SASS निर्देश शब्दकोश का रिवर्स इंजीनियरिंग, कर्नेल ऑडिट और GPU आर्किटेक्चर में पैटर्न पहचान।

रिपॉजिटरी देखेंवेबसाइट

SASS King logo

SASS किंग

नियंत्रित कर्नेल से लेकर प्रोडक्शन ऑडिट तक NVIDIA SASS का रिवर्स इंजीनियरिंग।

लेख 1 · लेख 2 · ज्ञानकोष · पैटर्न लाइब्रेरी · SM120 निर्देश शब्दकोश · एन्कोडिंग नोट्स · यहाँ से शुरू करें · प्रोजेक्ट संरचना · टेंसर-कोर अध्याय · योगदान

Architecture Status License

SASS किंग NVIDIA SASS का एक व्यवस्थित रिवर्स-इंजीनियरिंग प्रोजेक्ट है, जो संकलित CUDA बाइनरी के अंदर उत्सर्जित मूल GPU निर्देश सेट है। यह प्रोजेक्ट SM120 / SM120a उपभोक्ता Blackwell हार्डवेयर से शुरू होता है और समय के साथ एक पूर्ण क्रॉस-आर्किटेक्चर ISA और पैटर्न लाइब्रेरी की ओर बढ़ता है।

लक्ष्य व्यावहारिक है: एक कर्नेल इंजीनियर को SASS डंप खोलने, कंपाइलर पैटर्न पहचानने, प्रदर्शन-प्रासंगिक संरचनाओं की पहचान करने, और बाइनरी को स्रोत-स्तरीय अनुकूलन निर्णयों से जोड़ने में मदद करना।

प्रोजेक्ट ने अपनी प्रारंभिक चरण 3 पैटर्न लाइब्रेरी पूरी कर ली है: 29 पुन: प्रयोज्य SASS हस्ताक्षर अब patterns/ के अंतर्गत औपचारिक रूप में हैं, knowledge/FINDINGS.md को पूर्ण साक्ष्य श्रृंखला के रूप में रखा गया है। अगला बड़ा कदम चरण 4 है: वास्तविक प्रोडक्शन कर्नेल पर उन पैटर्न को लागू करना।

त्वरित नेविगेशन

रिपॉजिटरी एक साक्ष्य पाइपलाइन के रूप में व्यवस्थित है:

root@kitploit:~
corpus/      नियंत्रित कर्नेल और रॉ SASS साक्ष्य
knowledge/   प्रोजेक्ट-व्यापी निष्कर्ष, निर्देश नोट्स और एन्कोडिंग नोट्स
patterns/    पुन: प्रयोज्य चरण 3 ऑडिट हस्ताक्षर
production/  चरण 4 वास्तविक-कर्नेल ऑडिट

यह क्यों मौजूद है

भावना में समान अंतिम व्यापक सार्वजनिक SASS रिवर्स-इंजीनियरिंग कार्य 2018 में वोल्टा और ट्यूरिंग पर जिया एट अल. का था। एम्पीयर, हॉपर और ब्लैकवेल ने निर्देश मिश्रण को काफी हद तक बदल दिया है: एसिंक कॉपी पथ, टेंसर-कोर परिवार, मैट्रिक्स लोड/स्टोर निर्देश, स्पार्स और स्केल्ड MMA रूप, और नए यूनिफ़ॉर्म-रजिस्टर प्रवाह।

SASS किंग नियंत्रित माइक्रो-कर्नेल, रॉ SASS पढ़ना, रनटाइम प्रोब और प्रोडक्शन-कर्नेल ऑडिट को मिलाकर उस अंतर को भरता है।

वर्तमान स्थिति

चरण 3 प्रदेय

औपचारिक पैटर्न लाइब्रेरी चरण 3 का मुख्य आउटपुट है। यह अध्याय-स्थानीय साक्ष्य को पुन: प्रयोज्य ऑडिट हस्ताक्षर में बदलता है, ताकि एक ऑडिट हर बार पूर्ण शोध श्रृंखला को फिर से लिखने के बजाय एक नामित पैटर्न का हवाला दे सके।

चरण 3 को पूर्ण माना जाता है क्योंकि:

  • अध्याय 01-25 में पाए गए आवर्ती संरचनाओं को 29 नामित पैटर्न पृष्ठों में बढ़ावा दिया गया है;
  • प्रत्येक पैटर्न में सादा-अंग्रेजी स्पष्टीकरण, SASS हस्ताक्षर, वेरिएंट, एंटी-पैटर्न, खुले अंतर और आत्मविश्वास स्तर है;
  • दावा टैग knowledge/FINDINGS.md में स्रोत साक्ष्य से बंधे रहते हैं;
  • ऑडिट-सामने नेविगेशन अब patterns/README.md से शुरू होता है;
  • अनसुलझे आइटम को स्पष्ट रूप से अंतर के रूप में आगे बढ़ाया जाता है, पैटर्न टेक्स्ट के अंदर छिपाने के बजाय।

प्रत्येक पैटर्न पृष्ठ में शामिल हैं:

  • सादा-अंग्रेजी अर्थ;
  • SASS हस्ताक्षर;
  • देखे गए वेरिएंट;
  • व्याख्या सीमाएँ;
  • एंटी-पैटर्न;
  • खुले अंतर;
  • आत्मविश्वास स्तर।

ऑडिट-सामने सूचकांक के रूप में patterns/README.md का उपयोग करें। जब आपको किसी पैटर्न के पीछे लंबे शोध संदर्भ की आवश्यकता हो, तो knowledge/FINDINGS.md का उपयोग करें।

चरण 3 यह दावा नहीं करता कि प्रत्येक NVIDIA SASS व्यवहार डीकोड किया गया है। यह एक पुन: प्रयोज्य SM120 / SM120a पैटर्न परत स्थापित करता है जो मैन्युअल प्रोडक्शन ऑडिट शुरू करने के लिए पर्याप्त है। रनटाइम लेआउट डीकोड, पूर्ण नियंत्रण-कोड बिट प्लेसमेंट, स्वचालित क्यूबिन रिपोर्टिंग और क्रॉस-आर्किटेक्चर रीप्ले भविष्य के कार्य बने हुए हैं।

यहाँ से शुरू करें

  • प्रोजेक्ट में नए हैं: Start Here पढ़ें।
  • प्रोजेक्ट-व्यापी मानचित्र चाहिए: ज्ञानकोष सूचकांक पढ़ें।
  • वर्तमान निर्देश मानचित्र चाहिए: SM120 / SM120a पर SASS निर्देश पढ़ें।
  • एन्कोडिंग नोट्स चाहिए: एन्कोडिंग नोट्स पढ़ें।
  • कच्चा सत्य स्रोत चाहिए: निष्कर्ष पढ़ें।
  • पुन: प्रयोज्य ऑडिट हस्ताक्षर चाहिए: पैटर्न लाइब्रेरी पढ़ें।
  • टेंसर-कोर साक्ष्य चाहिए: टेंसर-कोर अध्याय से शुरू करें।
  • डंप या सुधार योगदान करना चाहते हैं: योगदान पढ़ें।
  • v0.1 सीमा चाहिए: रिलीज़ नोट्स पढ़ें।

सार्वजनिक लेख:

  • भाग 1 - प्रथम सिद्धांतों से NVIDIA SASS पढ़ना
  • भाग 2 - कंपाइलर का मन पढ़ना

पद्धति

नियंत्रित भिन्नता. दो कर्नेल ठीक एक चर में भिन्न होते हैं: डेटाटाइप, ऑपरेंड ऑर्डर, अनरोल फैक्टर, मेमोरी लेआउट, या संकलन लक्ष्य। SASS अंतर कंपाइलर निर्णय को अलग करता है।

सख्त दावा टैग. प्रत्येक तकनीकी दावा एक टैग का उपयोग करता है:

ऊपर-से-नीचे और नीचे-से-ऊपर एक साथ। माइक्रो-कर्नेल व्यक्तिगत निर्देशों और कंपाइलर निर्णयों को अलग करते हैं। प्रोडक्शन-जैसे कर्नेल दिखाते हैं कि वास्तविक कोड में कौन से पैटर्न मायने रखते हैं।

पैटर्न-प्रथम ऑडिट। एक प्रोडक्शन ऑडिट को दृश्य SASS हस्ताक्षर से मिलान करने और अपने आत्मविश्वास सीमाओं, एंटी-पैटर्न और खुले अंतर को स्थानांतरित करने के बाद ही एक औपचारिक PATTERN-NN पृष्ठ का हवाला देना चाहिए।

क्या शामिल है

पहला पास SM120 टेंसर-कोर और मेमोरी पाइपलाइन पर केंद्रित है:

  • HMMA, QMMA, OMMA
  • LDSM, STSM
  • LDGSTS, LDGDEPBAR, DEPBAR
  • LDG, STG, LDS, STS, REDG
  • BRA, EXIT, BSSY, ,

प्रोजेक्ट यह दिखावा नहीं करता कि ISA अभी भी पूर्ण है। सार्वजनिक शब्दकोश ट्रैक करता है कि क्या देखा और समझाया गया है; knowledge/encoding/ के तहत गहरे पृष्ठ उन परिवारों को ट्रैक करते हैं जिनमें मैचर-शैली दस्तावेज़ीकरण के लिए पर्याप्त साक्ष्य हैं।

SASS किंग बिट-स्तरीय SASS डिसअसेंबलरों से प्रतिस्पर्धा नहीं करता है। प्रोजेक्ट स्थानीय डंप को प्राथमिक साक्ष्य के रूप में उपयोग करता है और निर्देश क्षेत्रों, शेड्यूलिंग तालिकाओं, प्रेडिकेट्स और रजिस्टर ट्रैकिंग के लिए redplait/denvdis को क्रॉस-चेक के रूप में उपयोग कर सकता है। denvdis निम्न-स्तरीय एन्कोडिंग व्याख्याओं को मान्य कर सकता है; SASS किंग नियंत्रित-भिन्नता साक्ष्य, सिमैंटिक पैटर्न परत और प्रोडक्शन-ऑडिट व्याख्या का मालिक है।

रोडमैप

root@kitploit:~
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;

चरण 1 - शिक्षण कर्नेल

कर्नेल 01-12 बेसलाइन SASS अवधारणाओं को स्थापित करते हैं: FMA फ्यूजन, स्कोरबोर्ड व्यवहार, लूप लोअरिंग, शेयर्ड मेमोरी, ग्लोबल मेमोरी, वार्प प्राइमिटिव, स्लो-पथ गणित और स्थानीय-मेमोरी स्पिल।

चरण 2 - टेंसर-कोर और SM120 कवरेज

कर्नेल 13-25 वर्तमान SM120 टेंसर-कोर पथ को कवर करते हैं:

चरण 2.5 - denvdis क्रॉस-सत्यापन

प्रोडक्शन ऑडिट पैटर्न लाइब्रेरी पर निर्भर होने से पहले 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 आउटपुट सहायक साक्ष्य है, स्थानीय डंप अवलोकनों का प्रतिस्थापन नहीं।

चरण 3 - पैटर्न लाइब्रेरी

आवर्ती संरचनाओं को पुन: प्रयोज्य हस्ताक्षर में औपचारिक रूप दिया गया:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • श्रृंखलाबद्ध HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • वार्प कमी और क्रॉस-लेन सामूहिक
  • रजिस्टर-स्पिल हस्ताक्षर
  • स्केलर और यूनिफ़ॉर्म नियंत्रण-प्रवाह पैटर्न

प्रारंभिक चरण 3 लाइब्रेरी में patterns/ के अंतर्गत 29 पैटर्न पृष्ठ हैं। knowledge/FINDINGS.md शोध लॉग और सत्य स्रोत बना हुआ है; patterns/ ऑडिट-सामने प्रवेश बिंदु है।

चरण 3 प्रारंभिक लाइब्रेरी स्तर पर पूर्ण है। रनटाइम लेआउट डीकोड, पूर्ण नियंत्रण-कोड बिट प्लेसमेंट और क्रॉस-आर्किटेक्चर रीप्ले जैसे शेष आइटम अंतर या भविष्य के चरणों के रूप में ट्रैक किए जाते हैं, चरण 4 प्रोडक्शन ऑडिट शुरू करने में अवरोधक नहीं।

चरण 4 - प्रोडक्शन ऑडिट

FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad और संबंधित परियोजनाओं जैसी लाइब्रेरीज़ के वास्तविक कर्नेल पर पैटर्न लाइब्रेरी लागू करें। लक्ष्य एल्गोरिथमिक पैटर्न द्वारा प्रतिनिधित्वपूर्ण कवरेज है, प्रति कर्नेल एक मार्कडाउन फ़ाइल नहीं।

पहला चरण 4 प्रदेय एक मैन्युअल ऑडिट रिपोर्ट होना चाहिए जो:

  • एक वास्तविक कर्नेल को SASS क्षेत्रों में विभाजित करता है;
  • मिलान PATTERN-NN पृष्ठों का हवाला देता है;
  • प्रत्येक निष्कर्ष को आत्मविश्वास स्तर निर्दिष्ट करता है;
  • अस्पष्टीकृत क्षेत्रों को नए अंतर के रूप में रिकॉर्ड करता है;
  • जब तक कम से कम एक मैन्युअल रिपोर्ट स्थिर न हो जाए, ऑडिट उपकरण बनाने से बचता है।

चरण 5 - ऑडिट उपकरण

एक पाइपलाइन बनाएं जो एक क्यूबिन लेता है, ज्ञात पैटर्न का पता लगाता है, और अनुकूलन-उन्मुख रिपोर्ट जारी करता है।

चरण 6 - क्रॉस-आर्किटेक्चर

अतिरिक्त लक्ष्यों पर पद्धति को दोहराएं:

रिपॉजिटरी मैप

root@kitploit:~
.
├── 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 स्कोरबोर्ड, स्टॉल, दबाव और निर्भरता तीरों के लिए।
  • Nsight Compute प्रोफाइलिंग और स्टॉल विशेषता के लिए।
  • %clock माइक्रोबेंचमार्क निर्देश विलंबता जांच के लिए।
  • nvcc -Xptxas -v रजिस्टर और स्पिल मेटाडेटा के लिए।

संबंधित कार्य

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" विलंबता, थ्रूपुट और निर्भरता सत्यापन के पीछे अनुभवजन्य माइक्रोबेंचमार्किंग अनुशासन के लिए।
  • kuterdinel.com/nv_isa फ़ज़्ड NVIDIA ISA एन्कोडिंग कार्य के लिए, विशेष रूप से डिसअसेंबलर व्यवहार से मशीन-पढ़ने योग्य एन्कोडिंग नियम प्राप्त करने का विचार।
  • redplait/denvdis ऑपकोड तालिकाओं, बिट-स्तरीय डिसअसेंबली, एन्कोडिंग-फ़ील्ड निरीक्षण, शेड्यूलिंग विश्लेषण, रजिस्टर ट्रैकिंग और क्यूबिन हेरफेर के लिए। निकाले गए SM120 डेटा12 तालिकाओं का उपयोग निम्न-स्तरीय क्रॉस-चेक के रूप में किया जाता है जबकि स्थानीय डंप प्राथमिक साक्ष्य बने रहते हैं।
  • Redplait उपकरण और नोट्स: ced cubin editor, SASS disassembly Perl bindings, SASS latency analysis, और libcuda/nvasm_internal notes।
  • Huerta et al. 2025 कंपाइलर-निर्देशित शेड्यूलिंग, कंट्रोल कोड, डिपेंडेंसी काउंटर, रीयूज़ फ्लैग और यील्ड व्यवहार के रिवर्स-इंजीनियरिंग के लिए।
  • Yan et al. 2026 SASS के नीचे ड्राइवर-लेयर लॉन्च और पुशबफर विश्लेषण के लिए।
  • MaxAS और TuringAS पुराने NVIDIA आर्किटेक्चर के लिए पिछले सार्वजनिक SASS असेंबलर प्रयासों के रूप में।
  • NVIDIA CUDA Binary Utilities दस्तावेज़ीकरण आधिकारिक क्यूबिन, फैटबिन और डिसअसेंबली टूलिंग के लिए।

SASS किंग एल्गोरिथमिक पैटर्न परत पर काम करता है: संकलित कर्नेल कैसे संरचित होते हैं, इसे पहचानना और उन संरचनाओं को स्रोत-स्तरीय अनुकूलन निर्णयों से जोड़ना।

योगदान

योगदान का स्वागत है, विशेष रूप से:

  • रॉ SASS डंप जो यहाँ सीधे उपलब्ध नहीं हैं;
  • नियंत्रित कर्नेल अध्ययन जो एक कंपाइलर निर्णय को अलग करते हैं;
  • मौजूदा अवलोकनों में सुधार;
  • नए प्रोडक्शन-कर्नेल पैटर्न प्रस्ताव;
  • क्रॉस-आर्किटेक्चर तुलना।

अपेक्षित मेटाडेटा और लेखन मानक के लिए CONTRIBUTING.md देखें।

लेखक

फ्लोरियन मट्टाना। florianmattana.com

टूल डाउनलोड करें
यदि आप चाहते हैं...यहाँ से शुरू करेंफिर पढ़ें
10 मिनट में प्रोजेक्ट समझेंdocs/README.mddocs/START_HERE.md, फिर docs/PROJECT_STRUCTURE.md
साक्ष्य को पुन: उत्पन्न करेंcorpus/README.mdएक अध्याय conclusion*.md, फिर इसका .sass डंप
सत्य का स्रोत खोजेंknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
किसी नए डंप में एक पैटर्न पहचानेंpatterns/README.mdसंबंधित patterns/NN-*.md पृष्ठ
प्रोडक्शन ऑडिट शुरू करेंproduction/README.mdसंबंधित PATTERN-NN पृष्ठ और स्रोत साक्ष्य
कोई सुधार या डंप योगदान करेंCONTRIBUTING.mddocs/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-*
नियंत्रण प्रवाहविचलन/पुनर्मिलन, लूप बैक-एज, प्रेडिकेटेड एग्जिट, कोल्ड ट्रैप, स्थानीय CALLspatterns/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]स्पष्ट रूप से दस्तावेजीकृत खुला प्रश्न।
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • यूनिफ़ॉर्म-रजिस्टर प्रवाह: S2UR, 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 व्यवहार से अलग करता है।
    कर्नेलविषय
    13HMMA बेसलाइन, रजिस्टर आवंटन, संचायक श्रृंखलाबद्ध
    14QMMA FP8 / FP6 / FP4 बेसलाइन
    15संकीर्ण MMA वेरिएंट
    16FP4 पीक और ब्लॉक-स्केल्ड OMMA/QMMA
    17LDSM और मैट्रिक्स-लोड व्यवहार
    18पाइपलाइन्ड MMA टाइल और एसिंक कॉपी स्टेजिंग
    19स्पार्स MMA मेटाडेटा
    20नियंत्रण प्रवाह और बैक-एज डिटेक्शन
    21विचलन और पुनर्मिलन
    22STSM मैट्रिक्स-स्टोर व्यवहार
    23FP4 / FP6 फ़्रैगमेंट लेआउट प्रोब
    24प्रोडक्शन मिनी-GEMM ऑडिट
    25STSM एपिलॉग लेआउट और स्टोरबैक सिमैंटिक्स
    आर्कप्रतिनिधि GPUक्यों
    SM80A100डेटासेंटर Ampere बेसलाइन
    SM86RTX 3090उपभोक्ता Ampere कोर्पस
    SM89RTX 4090सामान्य उपभोक्ता अनुमान कार्ड
    SM90aH100TMA, WGMMA, वार्प स्पेशलाइज़ेशन, क्लस्टर
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090उपभोक्ता Blackwell प्रारंभिक बिंदु