Skip to content
KitploitKITPLOIT
도구블로그
제출
도구블로그
제출

해킹, 침투 테스트 및 사이버 보안 도구를 당신의 보안 무기고에!

Kitploit은 해킹, 사이버 보안 및 침투 테스트 도구 디렉토리입니다. 최신 프로젝트 업데이트를 발견하여 취약점을 찾고, 시스템을 분석하고, 테스트를 자동화하고, 보안을 강화하세요.

··피드·문의·개인정보·© 2026 Kitploit

도구 디렉토리

카테고리

모든 카테고리 보기
Loading categories
sass-king — NVIDIA SASS 명령어 사전의 리버스 엔지니어링, GPU 아키텍처 전반에 걸친 커널 감사 및 패턴 인식. | Kitploit
도구/GitHubGitHub/florianmattana/sass-king
Embedded Systems SecurityStatic AnalysisCode AnalysisReverse EngineeringHardware SecurityHardware & IoT SecurityBinary AnalysisPapers & ResearchLearning & EducationCurated ResourcesFirmware Analysis
317153개월 전Kitploit 검토 완료

인기

모두 보기 →

커뮤니티에서 가장 많이 사용되는 도구를 찾아보세요.

모든 도구 탐색

도구 컬렉션을 둘러보세요

모든 도구 보기 →
공유
GitHub
florianmattana/sass-king

sass-king

NVIDIA SASS 명령어 사전의 리버스 엔지니어링, GPU 아키텍처 전반에 걸친 커널 감사 및 패턴 인식.

저장소 보기웹사이트

SASS King 로고

SASS King

제어된 커널부터 프로덕션 감사까지 NVIDIA SASS를 리버스 엔지니어링합니다.

기사 1 · 기사 2 · 지식 기반 · 패턴 라이브러리 · SM120 명령어 용어집 · 인코딩 노트 · 여기서 시작 · 프로젝트 구조 · 텐서 코어 챕터 · 기여하기

Architecture Status License

SASS King은 NVIDIA SASS(컴파일된 CUDA 바이너리 내에서 방출되는 네이티브 GPU 명령어 집합)에 대한 체계적인 리버스 엔지니어링 프로젝트입니다. 이 프로젝트는 SM120 / SM120a 소비자 Blackwell 하드웨어에서 시작하여 시간이 지남에 따라 완전한 크로스 아키텍처 ISA 및 패턴 라이브러리로 확장됩니다.

목표는 실용적입니다: 커널 엔지니어가 SASS 덤프를 열고, 컴파일러 패턴을 인식하고, 성능 관련 구조를 식별하고, 바이너리를 다시 소스 수준 최적화 결정에 연결할 수 있도록 돕습니다.

프로젝트는 초기 Phase 3 패턴 라이브러리를 완료했습니다: 29개의 재사용 가능한 SASS 시그니처가 patterns/ 아래에 공식화되었으며, knowledge/FINDINGS.md는 전체 증거 추적으로 유지됩니다. 다음 주요 단계는 Phase 4: 이러한 패턴을 실제 프로덕션 커널에 적용하는 것입니다.

빠른 탐색

저장소는 증거 파이프라인으로 구성됩니다:

root@kitploit:~
corpus/      제어된 커널 및 원시 SASS 증거
knowledge/   프로젝트 전체 결과, 명령어 노트 및 인코딩 노트
patterns/    재사용 가능한 Phase 3 감사 시그니처
production/  Phase 4 실제 커널 감사

존재 이유

비슷한 정신의 마지막 광범위한 공개 SASS 리버스 엔지니어링 작업은 2018년 Volta 및 Turing에 대한 Jia 등의 연구였습니다. Ampere, Hopper 및 Blackwell은 명령어 구성을 상당히 변경했습니다: 비동기 복사 경로, 텐서 코어 제품군, 행렬 로드/저장 명령어, 희소 및 스케일링된 MMA 형식, 새로운 균일 레지스터 흐름.

SASS King은 제어된 마이크로 커널, 원시 SASS 읽기, 런타임 프로브 및 프로덕션 커널 감사를 결합하여 이 격차를 메웁니다.

현재 상태

Phase 3 인도물

공식 패턴 라이브러리는 Phase 3의 주요 결과물입니다. 챕터별 증거를 재사용 가능한 감사 시그니처로 변환하여, 감사 시 전체 연구 추적을 매번 다시 작성하는 대신 명명된 패턴을 인용할 수 있습니다.

Phase 3은 다음 이유로 완료된 것으로 간주됩니다:

  • 챕터 01-25에서 발견된 반복 구조가 29개의 명명된 패턴 페이지로 승격되었습니다.
  • 모든 패턴에는 일반 영어 설명, SASS 시그니처, 변형, 안티 패턴, 열린 격차 및 신뢰 수준이 있습니다.
  • 주장 태그는 knowledge/FINDINGS.md의 소스 증거에 범위가 지정됩니다.
  • 감사 지향 탐색은 이제 patterns/README.md에서 시작됩니다.
  • 해결되지 않은 항목은 패턴 텍스트 내에 숨겨지지 않고 격차로 명시적으로 전달됩니다.

각 패턴 페이지에는 다음이 포함됩니다:

  • 일반 영어 의미;
  • SASS 시그니처;
  • 관찰된 변형;
  • 해석 경계;
  • 안티 패턴;
  • 열린 격차;
  • 신뢰 수준.

감사 지향 색인으로 patterns/README.md를 사용하세요. 패턴 뒤에 있는 더 긴 연구 맥락이 필요하면 knowledge/FINDINGS.md를 사용하세요.

Phase 3은 모든 NVIDIA SASS 동작이 디코딩되었다고 주장하지 않습니다. 수동 프로덕션 감사를 시작하기에 충분한 재사용 가능한 SM120 / SM120a 패턴 레이어를 설정합니다. 런타임 레이아웃 디코드, 전체 제어 코드 비트 배치, 자동화된 cubin 보고 및 크로스 아키텍처 재생은 향후 작업으로 남습니다.

여기서 시작

  • 프로젝트가 새롭다면: 여기서 시작을 읽으세요.
  • 프로젝트 전체 지도를 원한다면: 지식 기반 색인을 읽으세요.
  • 현재 명령어 맵을 원한다면: SM120 / SM120a의 SASS 명령어를 읽으세요.
  • 인코딩 노트를 원한다면: 인코딩 노트를 읽으세요.
  • 원시 진리의 원천을 원한다면: 결과를 읽으세요.
  • 재사용 가능한 감사 시그니처를 원한다면: 패턴 라이브러리를 읽으세요.
  • 텐서 코어 증거를 원한다면: 텐서 코어 챕터로 시작하세요.
  • 덤프나 수정 사항을 기여하려면: 기여하기를 읽으세요.
  • v0.1 경계를 원한다면: 릴리스 노트를 읽으세요.

공개 글:

  • Part 1 - 첫 번째 원칙에서 NVIDIA SASS 읽기
  • Part 2 - 컴파일러의 마음 읽기

방법론

제어된 변형. 두 커널은 정확히 하나의 변수만 다릅니다: dtype, 피연산자 순서, 언롤 팩터, 메모리 레이아웃 또는 컴파일 대상. SASS diff는 컴파일러 결정을 분리합니다.

엄격한 주장 태그. 모든 기술적 주장은 태그를 사용합니다:

상향식과 하향식을 함께. 마이크로 커널은 개별 명령어와 컴파일러 결정을 분리합니다. 프로덕션 유사 커널은 실제 코드에서 어떤 패턴이 중요한지 보여줍니다.

패턴 우선 감사. 프로덕션 감사는 가시적인 SASS 시그니처와 일치하고 신뢰 한계, 안티 패턴 및 열린 격차를 이어받은 후에만 공식 PATTERN-NN 페이지를 인용해야 합니다.

포함되는 내용

첫 번째 패스는 SM120 텐서 코어 및 메모리 파이프라인에 초점을 맞춥니다:

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

프로젝트는 ISA가 아직 완전하지 않다고 주장하지 않습니다. 공개 용어집은 관찰되고 설명된 내용을 추적합니다. knowledge/encoding/ 아래의 더 깊은 페이지는 매처 스타일 문서화에 충분한 증거가 있는 제품군을 추적합니다.

SASS King은 비트 수준 SASS 디스어셈블러와 경쟁하지 않습니다. 프로젝트는 로컬 덤프를 기본 증거로 사용하며 명령어 필드, 스케줄링 테이블, 조건자 및 레지스터 추적에 대한 교차 확인을 위해 redplait/denvdis를 사용할 수 있습니다. denvdis는 저수준 인코딩 해석을 검증할 수 있습니다. SASS King은 제어 변형 증거, 의미 패턴 레이어 및 프로덕션 감사 해석을 소유합니다.

로드맵

root@kitploit:~
flowchart LR
    P1["Phase 1<br/>Teaching kernels<br/>01-12"] --> P2["Phase 2<br/>SM120 tensor-core corpus<br/>13-25"]
    P2 --> P25["Phase 2.5<br/>denvdis cross-validation<br/>bit-level backend"]
    P25 --> P3["Phase 3<br/>Pattern library<br/>compiler signatures"]
    P3 --> P4["Phase 4<br/>Production audits<br/>real kernels"]
    P4 --> P5["Phase 5<br/>Audit tool<br/>cubin reports"]
    P5 --> P6["Phase 6<br/>Cross-architecture replay<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;

Phase 1 - 교육 커널

커널 01-12는 기본 SASS 개념을 수립합니다: FMA 융합, 스코어보드 동작, 루프 로우어링, 공유 메모리, 글로벌 메모리, 워프 프리미티브, 느린 경로 수학 및 로컬 메모리 스필.

Phase 2 - 텐서 코어 및 SM120 커버리지

커널 13-25는 현재 SM120 텐서 코어 경로를 다룹니다:

Phase 2.5 - denvdis 교차 검증

프로덕션 감사가 패턴 라이브러리에 의존하기 전에 redplait/denvdis를 SM120 / SM120a의 비트 수준 교차 확인 백엔드로 검증합니다. 패스는 대표적인 로컬 cubin 또는 HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR 및 분기 마커를 다루는 덤프에서 nvd -O, nvd -S, nvd -p 및 유용한 경우 nvd -T를 실행합니다.

출력은 knowledge/DENVDIS_INTEGRATION.md입니다: 제품군별 denvdis 인식 상태, 수정자 커버리지, 노출된 제어 코드 필드 및 SASS King 조치의 사실적 호환성 표입니다. denvdis 출력은 로컬 덤프 관찰을 대체하지 않는 지원 증거입니다.

Phase 3 - 패턴 라이브러리

반복되는 구조를 재사용 가능한 시그니처로 공식화했습니다:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • 체이닝된 HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • 워프 리덕션 및 크로스 레인 컬렉티브
  • 레지스터 스필 시그니처
  • 스칼라 및 균일 제어 흐름 패턴

초기 Phase 3 라이브러리에는 patterns/ 아래 29개의 패턴 페이지가 포함되어 있습니다. knowledge/FINDINGS.md는 연구 로그 및 진리의 원천으로 남아 있습니다. patterns/는 감사 지향 진입점입니다.

Phase 3은 초기 라이브러리 수준에서 완료되었습니다. 런타임 레이아웃 디코드, 전체 제어 코드 비트 배치 및 크로스 아키텍처 재생과 같은 나머지 항목은 Phase 4 프로덕션 감사 시작을 차단하는 것이 아니라 격차 또는 향후 단계로 추적됩니다.

Phase 4 - 프로덕션 감사

패턴 라이브러리를 FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad 및 관련 프로젝트와 같은 라이브러리의 실제 커널에 적용합니다. 목표는 커널당 하나의 마크다운 파일이 아니라 알고리즘 패턴별로 대표적인 커버리지입니다.

첫 번째 Phase 4 인도물은 다음과 같은 수동 감사 보고서여야 합니다:

  • 하나의 실제 커널을 SASS 영역으로 세분화;
  • 일치하는 PATTERN-NN 페이지 인용;
  • 각 결론에 신뢰 수준 할당;
  • 설명되지 않은 영역을 새로운 격차로 기록;
  • 최소한 하나의 수동 보고서가 안정될 때까지 감사 도구 구축을 피합니다.

Phase 5 - 감사 도구

cubin을 가져와 알려진 패턴을 감지하고 최적화 지향 보고서를 생성하는 파이프라인을 구축합니다.

Phase 6 - 크로스 아키텍처

추가 대상에 방법론을 재적용:

저장소 맵

root@kitploit:~
.
├── corpus/                                # 제어된 커널, 덤프 및 챕터 글
│   ├── basics/                            # 커널 01-08: 스칼라/벡터 및 메모리 기본
│   ├── warp_collectives/                  # 커널 09-10: shuffle, vote, reduction
│   ├── math_and_spills/                   # 커널 11-12: 느린 경로 및 스필
│   └── tensor_cores/                      # 커널 13-25: 텐서 코어 연구
├── knowledge/                              # 결과, 용어집, 인코딩 노트
│   ├── FINDINGS.md
│   ├── SASS_INSTRUCTIONS_SM120.md
│   └── encoding/
├── patterns/                               # 공식 Phase 3 패턴 라이브러리
├── production/                             # Phase 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 - Opcode 테이블, 비트 수준 디스어셈블리, 인코딩 필드 검사, 스케줄링 분석, 레지스터 추적 및 cubin 조작용. 추출된 SM120 data12 테이블은 로컬 덤프가 기본 증거로 남아 있는 동안 저수준 교차 확인으로 사용됩니다.
  • 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 문서 - 공식 cubin, fatbin 및 디스어셈블리 도구용.

SASS King은 알고리즘 패턴 레이어에서 작동합니다: 컴파일된 커널이 어떻게 구조화되어 있는지 인식하고 이러한 구조를 소스 수준 최적화 결정에 연결합니다.

기여하기

기여를 환영하며, 특히 다음을 환영합니다:

  • 여기서 직접 사용할 수 없는 하드웨어의 원시 SASS 덤프;
  • 하나의 컴파일러 결정을 분리하는 제어된 커널 연구;
  • 기존 관찰에 대한 수정;
  • 새로운 프로덕션 커널 패턴 제안;
  • 크로스 아키텍처 비교.

예상되는 메타데이터 및 작성 표준은 CONTRIBUTING.md를 참조하세요.

저자

Florian Mattana. 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
패턴 라이브러리초기 Phase 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-*
제어 흐름분기/재수렴, 루프 백에지, 조건부 종료, 콜드 트랩, 로컬 CALLpatterns/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-*
워프 컬렉티브워프 리덕션, shuffle/vote/match/sync 프리미티브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-to-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. 감사 도구계획됨cubin-to-보고서 파이프라인패턴 레이어를 스크립팅 가능하고 반복 가능하게 만듭니다.
    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 시작점