Thrust → CUDA → PTX → SASS 🏋️‍♂️🏋️‍♀️ (#25)

* Add: Thrust, CUB, CUDA sorting

This is a draft. It still lacks manual
timing and async scheduling.

* Add: Thrust, CUB, CUDA sorting

This is a draft. It still lacks manual
timing and async scheduling.

* Make: Options for CUDA & TBB in CMake

* Make: Switch to CUDA Toolkit for GPU libs

* Fix: Ranges require `constexpr` on NVCC

* Make: Upgrade `fmt` for NVCC builds

fmtlib/fmt#4297

* Fix: NVCC compilation issues

* Make: Silence NVCC warnings

* Add: Sorting with `thrust` and `cub`

* Add: PTX and `.cuh` kernels

* Make: Don't compile PTX

* Add: Using CUDA Driver API to JIT `.ptx`

* Add: Precompiled CUDA C++ kernels

* Add: cuBLAS benchmarks

* Fix: Compiling `cuBLAS` calls

* Fix: Avoid optimizing-out SASS code

Unless we put an impossible condition with
a `wmma::store_matrix_sync` the result of
fragment multiplication is optimized out.

* Add: Tensor Core intrinsic benchmarks

Targeting `f16`, `bf16`, `tf16`, `f32`, `f64`
on Volta, Turing, and Ampere.

* Make: Build CUDA for multiple platforms

Currently covering Volta, Turing, Ampere,
Ada Lovelace, and Hopper.

* Add: Binary BMMA kernels for GPU

XOR variant for Turing+.
AND variant for Ampere+.

* Docs: Introduce Warp-Group-MMA on Hopper

* Fix: Working PTX kernel

* Fix: Lower PTX version for JIT

* Fix: Use `f16` MMA

* Make: Drop OpenBLAS

Author: ashvardanian

.vscode/settings.json

@@ -6,20 +6,26 @@
     "Andreas",
     "ashvardanian",
     "asynchrony",
+    "bfloat",
     "bioinformatics",
     "BLAS",
+    "bmma",
     "Boccara",
     "bootcamps",
     "Byrne",
     "Cawley",
     "cblas",
+    "CCCL",
     "clflush",
     "colsb",
     "consteval",
+    "constexpr",
     "coro",
     "cplusplus",
     "cppcoro",
+    "cstddef",
     "CTRE",
+    "cublas",
     "CUDA",
     "denormal",
     "DOTPROD",
@@ -28,20 +34,26 @@
     "Eron",
     "excerise",
     "fconcepts",
+    "fdividef",
     "Fedor",
     "Fugaku",
     "Giga",
     "Goodput",
     "GOPS",
+    "GPGPU",
     "grandkids",
+    "GTEST",
     "Hana",
+    "HKSTU",
+    "HMMA",
     "Ibireme",
     "JeanHeyd",
     "jemalloc",
     "Kulukundis",
     "Lelbach",
     "Lemire",
     "Lib",
+    "LIBPFM",
     "libunifex",
     "Lohmann",
     "Maclaurin",
@@ -54,19 +66,27 @@
     "Niebler",
     "Niels",
     "nlohmann",
+    "noexcept",
+    "nullptr",
     "NVCC",
+    "nvcuda",
     "openblas",
     "openmp",
     "Ormrod",
     "Peta",
     "Pikus",
     "pmf",
+    "POPCOUNT",
     "popcountll",
+    "Pranjal",
     "prefetcher",
     "pthread",
+    "PTXAS",
     "RDMA",
     "reorderable",
+    "Shankhdhar",
     "simdjson",
+    "sinf",
     "SLEEF",
     "spdlog",
     "STREQUAL",
@@ -86,7 +106,11 @@
     "Vardanian",
     "vfmadd",
     "VNNI",
+    "VPCLMULQDQ",
     "Weis",
+    "WGMMA",
+    "wmma",
+    "Worklog",
     "XCOMP",
     "XFEATURE",
     "XTILE",
@@ -97,7 +121,10 @@
     "Zverovich"
   ],
   "files.associations": {
+    "*.evaluator": "cpp",
+    "*.ipp": "cpp",
     "*.tcc": "cpp",
+    "*.traits": "cpp",
     "__bit_reference": "cpp",
     "__bits": "cpp",
     "__config": "cpp",
@@ -114,12 +141,16 @@
     "__tree": "cpp",
     "__tuple": "cpp",
     "__verbose_abort": "cpp",
+    "adolcforward": "cpp",
     "algorithm": "cpp",
+    "alignedvector3": "cpp",
     "any": "cpp",
     "array": "cpp",
     "atomic": "cpp",
+    "autodiff": "cpp",
     "bit": "cpp",
     "bitset": "cpp",
+    "bvh": "cpp",
     "cctype": "cpp",
     "cfenv": "cpp",
     "charconv": "cpp",
@@ -130,6 +161,7 @@
     "complex": "cpp",
     "concepts": "cpp",
     "condition_variable": "cpp",
+    "core": "cpp",
     "coroutine": "cpp",
     "cstdarg": "cpp",
     "cstddef": "cpp",
@@ -141,30 +173,38 @@
     "cwchar": "cpp",
     "cwctype": "cpp",
     "deque": "cpp",
+    "eulerangles": "cpp",
     "exception": "cpp",
     "execution": "cpp",
     "expected": "cpp",
+    "fft": "cpp",
     "format": "cpp",
     "forward_list": "cpp",
     "fstream": "cpp",
     "functional": "cpp",
     "initializer_list": "cpp",
+    "inplace_vector": "cpp",
     "iomanip": "cpp",
     "ios": "cpp",
     "iosfwd": "cpp",
     "iostream": "cpp",
     "istream": "cpp",
     "iterator": "cpp",
+    "kroneckerproduct": "cpp",
     "limits": "cpp",
     "list": "cpp",
     "locale": "cpp",
     "map": "cpp",
     "memory": "cpp",
     "memory_resource": "cpp",
+    "mprealsupport": "cpp",
     "mutex": "cpp",
     "new": "cpp",
+    "nnls": "cpp",
     "numbers": "cpp",
     "numeric": "cpp",
+    "numericaldiff": "cpp",
+    "openglsupport": "cpp",
     "optional": "cpp",
     "ostream": "cpp",
     "random": "cpp",
@@ -176,12 +216,15 @@
     "shared_mutex": "cpp",
     "source_location": "cpp",
     "span": "cpp",
+    "specialfunctions": "cpp",
+    "splines": "cpp",
     "sstream": "cpp",
     "stdexcept": "cpp",
     "stop_token": "cpp",
     "streambuf": "cpp",
     "string": "cpp",
     "string_view": "cpp",
+    "superlusupport": "cpp",
     "system_error": "cpp",
     "text_encoding": "cpp",
     "thread": "cpp",
@@ -194,22 +237,16 @@
     "valarray": "cpp",
     "variant": "cpp",
     "vector": "cpp",
-    "core": "cpp",
-    "superlusupport": "cpp",
-    "*.evaluator": "cpp",
-    "*.traits": "cpp",
-    "adolcforward": "cpp",
-    "alignedvector3": "cpp",
-    "autodiff": "cpp",
-    "bvh": "cpp",
-    "eulerangles": "cpp",
-    "fft": "cpp",
-    "kroneckerproduct": "cpp",
-    "mprealsupport": "cpp",
-    "nnls": "cpp",
-    "numericaldiff": "cpp",
-    "openglsupport": "cpp",
-    "specialfunctions": "cpp",
-    "splines": "cpp"
+    "version": "cpp",
+    "strstream": "cpp",
+    "cinttypes": "cpp",
+    "codecvt": "cpp",
+    "csetjmp": "cpp",
+    "csignal": "cpp",
+    "hash_map": "cpp",
+    "future": "cpp",
+    "scoped_allocator": "cpp",
+    "typeindex": "cpp",
+    "*.inc": "cpp"
   }
 }