Über mich
Ich schreibe performancekritischen Systemcode und messe ihn an der besten Referenz, die ich bekommen kann, in der Regel die Herstellerbibliothek oder das Limit der Hardware. Jedes Projekt hier veröffentlicht seine Zahlen und das Messgerüst, das sie erzeugt hat, denn ein Speedup, den niemand reproduzieren kann, ist eine Behauptung und kein Ergebnis. Begonnen habe ich im Maschinenbau, wo ich physische Systeme gebaut habe. Ins Computational Engineering bin ich gewechselt, um die Software zu schreiben, die solche Systeme simuliert und beschleunigt.
Ausgewählte Projekte
-
CUDA-Kernel-Library
Eine GEMM-Staffel von einem naiven Kernel mit 1.908 GFLOP/s bis zu 54.985 GFLOP/s mit mma.sync, ldmatrix und geswizzeltem Shared Memory, also 90 % von cuBLAS in FP16 auf derselben GPU, bei 83 % des Speed-of-Light-Limits der Tensor-Pipe. Jede Stufe ist ein eigener Kernel, dessen Profiler-Lauf mit eingecheckt ist.
-
MLIR-Backend-for-a-Simulated-Edge-NPU
Ein Deep-Learning-Compiler-Backend, das ONNX über eigene MLIR-Dialekte in einen binären Befehlsstrom für eine simulierte NPU übersetzt. Bei LeNet senkt -O2 die Instruktionszahl von 91 auf 70, die simulierten Zyklen von 23.421 auf 12.710 und den DRAM-Traffic von 339 KB auf 177 KB; die Ausgabe stimmt bis auf 3e-8 mit onnxruntime überein.
-
Parallel-Numerical-Library
Eine C++23-Bibliothek, deren iterative Löser unverändert auf seriellen, OpenMP-, Thread-, MPI- und CUDA-Backends laufen und dabei bitidentische Ergebnisse liefern. Gemessen an der jeweils eigenen STREAM-Bandbreite des Geräts, 62,3 GiB/s auf der CPU und 549,7 GiB/s auf der GPU, erreicht Red-Black-Gauß-Seidel 45,7 % beziehungsweise 47,1 %. 440 Konfigurationen, in 23 Minuten vollständig reproduzierbar.
-
Particle-Physics-Simulation
Eine CUDA-Engine für N-Body-Gravitation und SPH-Fluide mit 17,47 TFLOP/s in FP32, wobei Shared-Memory-Tiling bei 200.000 Partikeln 1,93x schneller ist als der naive Kernel. Die Korrektheit sichern zehn Tests ab, darunter eine Energiedrift von 1,33e-05 über 10.000 Schritte.
-
Ring-AllReduce
Ein Ring-Allreduce, allein aus MPI-Punkt-zu-Punkt-Primitiven aufgebaut und anschließend an das Alpha-Beta-Kostenmodell angepasst: 0,250 us pro Nachricht, 305,4 ps pro Byte, R-Quadrat 0,981 über 375 Punkte. Die Dokumentation benennt ausdrücklich, wo das Verfahren gegen das Hersteller-Collective verliert und warum.
-
Low-Power-Async-CDC-FIFO-IP
Ein parametrierbarer asynchroner FIFO für Clock-Domain-Crossing in SystemVerilog, mit Gray-Code-Zeigern und Low-Power-Intent, in SymbiYosys beschränkt bewiesen und über 60 Konfigurationen hinweg auf Fläche, Leistung und Timing vermessen.
Weitere Repositories, darunter ein Out-of-Tree-LLVM-Pass-Framework, CPU-Mikrobenchmarks, GPU-Roofline-Profiling, MPI-Collective-Benchmarks und Werkzeuge für ML-Experimente, finden sich auf meinem GitHub-Profil.
Kenntnisse
- Programmiersprachen
- C++, Python, C, CUDA, SystemVerilog, Assembly (x86-64, AArch64), MATLAB
- GPU und HPC
- CUDA-Kernel-Optimierung, Tensor-Core-Programmierung (mma.sync, WMMA), OpenMP, MPI, pthreads und C++-Threads, SIMD (AVX2, NEON), Nsight-Profiling, Roofline- und Speed-of-Light-Analyse
- Compiler und Hardware
- MLIR, LLVM, ONNX-Lowering, RTL-Design und formale Verifikation (SymbiYosys), Yosys- und OpenSTA-Flows
- ML und Frameworks
- PyTorch, torch.distributed, NumPy, pandas, Experiment-Tracking und Trainingsdiagnostik
- Werkzeuge
- Git, CMake, Docker, Linux und WSL2, GitHub Actions, LaTeX, ANSYS, Siemens NX
Ausbildung und Erfahrung
- M.Sc. Computational Engineering, Ruhr-Universität Bochum, seit 2024, aktuelle Durchschnittsnote 1,9
- B.Eng. Maschinenbau, Federal University of Technology Akure, 2017 bis 2023, Abschlussnote 1,1
- Softwareentwickler, Whitelegg Nigeria Limited, Lagos, 2023 bis 2024
- Praktikant Data Engineering, Monitoring-Automatisierung, Chikki Foods Industries, Lagos, 2021