v0.7: Networking in POSIX vs. io_uring 💍

To showcase the differences between different IO approaches, this release brings a batch-asynchronous echo server implementation on top of UDP, measuring the packet drop frequency, throughput, and latency for:

• ASIO
• POSIX
• io_uring

The numbers currently look like:

Running build_release/less_slow
Run on (6 X 4000.4 MHz CPU s)
CPU Caches:
  L1 Data 48 KiB (x6)
  L1 Instruction 32 KiB (x6)
  L2 Unified 2048 KiB (x6)
  L3 Unified 327680 KiB (x1)
Load Average: 0.93, 0.52, 0.47
----------------------------------------------------------------------------------------------------------
Benchmark                                                Time             CPU   Iterations UserCounters...
----------------------------------------------------------------------------------------------------------
rpc_libc/loopback/min_time:2.000/manual_time          5514 us         2298 us          509 bytes_per_second=45.3389Mi/s drop,%=0 items_per_second=46.427k/s max_packet_latency,us=55 mean_batch_latency,us=5.51403k mean_packet_latency,us=21.5392
rpc_uring55/loopback/min_time:2.000/manual_time       1630 us         1591 us         1727 bytes_per_second=153.366Mi/s drop,%=0 items_per_second=157.046k/s max_packet_latency,us=1.822k mean_batch_latency,us=1.63009k mean_packet_latency,us=6.36754
rpc_asio/loopback/min_time:2.000/manual_time         89058 us          878 us           28 bytes_per_second=2.80717Mi/s drop,%=12.9325 items_per_second=2.87454k/s max_packet_latency,us=916 mean_batch_latency,us=89.0576k mean_packet_latency,us=399.553

The current example only uses the most basic io_uring features available with Linux kernel 5.5. In the next iterations (#30), we should extend it with the following functionality:

• IORING_REGISTER_BUFFERS - since 5.1
• IORING_RECV_MULTISHOT or io_uring_prep_recvmsg_multishot - since 6.0
• IORING_OP_SEND_ZC or io_uring_prep_sendmsg_zc - since 6.0
• IORING_SETUP_SQPOLL - with IORING_FEAT_SQPOLL_NONFIXED after 5.11
• IORING_SETUP_SUBMIT_ALL - since 5.18
• IORING_SETUP_COOP_TASKRUN - since 5.19
• IORING_SETUP_SINGLE_ISSUER - since 6.0

Feel free to join the development 🤗

Minor

• Add: io_uring variant for kernel 6.0 (ce73aa3)
• Add: io_uring draft (ec28b57)
• Add: External route networking (a2a8c9e)
• Add: POSIX echo implementation (3cce3b9)
• Add: ASIO "echo" server/client ping-pong (08d3326)

Patch

• Fix: Depend io_uring compilation on kernel version (70c53f6)
• Improve: IOSQE_FIXED_FILE for kernel 6.0+ (f7f7693)
• Improve: ASIO benchmarks (6be216a)
• Docs: Refactor spell-checks (24706a7)
• Make: Order spell-checks (1358a69)
• Docs: Recommend OpenBLAS (035e388)
• Improve: Avoid std::format in io_uring (1857a82)
• Fix: ARCH_ENABLE_TAGGED_ADDR needs Linux 6.2+ (3993b0c)
• Fix: Missing openblas_set_num_threads (3cab87d)
• Docs: Instal libBLAS (7629609)
• Improve: SO_ZEROCOPY (fd4c9e2)
• Improve: Retrofit registering buffers in 5.5 (95de751)
• Make: RelWithDebInfo flags (2838fd5)
• Improve: Code styling on Windows (c9238a1)
• Fix: Avoid in-place increment (2c25b4d)
• Make: Disable CUDA by default (94879fd)
• Make: Matching VERSION in CMake (b4dc186)
• Improve: Detect Linux version (f3e91fa)
• Improve: physical_cores for Windows refactor (0eb985c)
• Docs: Future io_uring tasks (53c4ca6)
• Improve: io_uring optional timeouts (f933582)
• Make: Revert to default BLAS (b3e13dd)
• Improve: io_uring server logic (3dfe612)
• Fix: liburing example (a2a9d6c)
• Improve: Reuse benchmarking logic (cae4175)
• Improve: Manual IO timing (fc60bfd)
• Make: Switch to PkgConfig for liburing (b4e50ad)
• Fix: Compiling asio example (b041392)
• Make: Tag dependencies, where possible (8f2e985)
• Improve: Batching client/server requests (955be1d)
• Make: liburing & asio deps (7256f98)

Less Slow v0.6: Thrust → CUDA → PTX → SASS 🏋️‍♂️🏋️‍♀️

It's almost impossible to imagine modern High-Performance Computing without GPUs. Yet, there are surprisingly few "full stack" demos out there for folks wanting to build intuition around CUDA C++, PTX Intermediate Representations, SASS Assembly, and higher-level libraries like Thrust, CUB, or the various cuBLAS flavors. This new release of Less Slow covers all of those! 🥳

Tensor Cores

The main highlight is an in-depth look at Tensor Core designs, from their extensive type system to the complexity of tile shapes—notoriously under-documented and confusing areas. These capabilities differ across Volta, Turing, Ampere, Ada, and Hopper GPUs, mapping to different PTX intrinsics (like wmma, binary bmma, or warp-group wgmma) and culminating in yet another shape at the SASS level with instructions such as multiple HMMA.884.F32.F32.STEPx instructions for each wmma.mma.sync.aligned.row.col.m16n16k16.f32.f32 intrinsic on Volta. And if you believe that instruction is long... be warned 😅

__global__ void tops_f16f16_sm70tc_16x16x16_1024unroll_cuda_kernel() {
    using namespace nvcuda;
    wmma::fragment<wmma::matrix_a, 16,16,16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16,16,16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16,16,16, half> c_frag;
    for (int i = 0; i < 1024; ++i)
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
}

$ cuobjdump -sass less_slow_from_cu.cubin | grep -i mma
# e.g. HMMA.884.F32.F32.STEP2 ...

This indicates the 8×8×4 shape actually used by the hardware on Volta.

PTX vs SASS

I've also hand-written PTX kernels, that may look like:

.visible .entry tops_f16f16_sm70tc_16x16x16_1024loop_ptx_kernel()
{
  // ...
  loop_start:
    // A single wmma instruction
    wmma.mma.sync.aligned.row.col.m16n16k16.f16.f16
      { %f0, %f1, %f2, %f3 }, // output accumulators
      { %f4, ... },          // A
      { %f12, ... },         // B
      { %f0, %f1, %f2, %f3 }; // input accumulators
    // ...
  bra loop_start;
}

Using the provided scripts, you can see for yourself just how different manually written vs. machine-generated PTX can be and how to invoke kernels directly from C++ in various ways — whether through the CUDA Runtime API or the CUDA Driver API — loading and JIT-compiling bits of PTX on the fly!

cuInit(0);
CUdevice dev; cuDeviceGet(&dev, 0);
CUcontext ctx; cuCtxCreate(&ctx, 0, dev);
CUmodule mod; cuModuleLoad(&mod, "less_slow.ptx");
CUfunction fun; cuModuleGetFunction(&fun, mod, "tops_f16f16_sm70tc_16x16x16_1024loop_ptx_kernel");

void* args[] = { /* kernel parameters here */ };
cuLaunchKernel(fun,
               1, 1, 1,  // gridDim
               256, 1, 1,// blockDim
               0, nullptr, args, nullptr);
cuCtxSynchronize();
cuModuleUnload(mod);
cuCtxDestroy(ctx);

cuBLAS on Practice

I've also included theoretical throughput benchmarks alongside real matrix multiplications via cuBLAS in case you want to compare actual performance to the raw theoretical numbers. One important observation here may be the lack of low-resolution numeric types:

if constexpr (std::is_same_v<scalar_type_, float>) {
    scalar_type_ alpha = 1, beta = 0;
    cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n, &alpha, a.begin(), lda, b.begin(), ldb, &beta, c.begin(), ldc);
} else if constexpr (std::is_same_v<scalar_type_, double>) {
    scalar_type_ alpha = 1, beta = 0;
    cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n, &alpha, a.begin(), lda, b.begin(), ldb, &beta, c.begin(), ldc);
} else if constexpr (std::is_same_v<scalar_type_, __half>) {
    scalar_type_ alpha = 1, beta = 0;
    cublasHgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n, &alpha, a.begin(), lda, b.begin(), ldb, &beta, c.begin(), ldc);
} else if constexpr (std::is_same_v<scalar_type_, int8_t>) {
    int32_t alpha_int = 1, beta_int = 0;
    cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n, &alpha_int, a.begin(), CUDA_R_8I, lda, b.begin(), CUDA_R_8I, ldb, &beta_int, c.begin(), CUDA_R_32I, ldc, CUDA_R_32I, CUBLAS_GEMM_DEFAULT);
}

Even integer kernels have a different signature, requiring $Alpha$ and $Beta$ to match the accumulator type, rather than the inputs. Very few libraries have adaptations for binary matrices and or sub-byte representations.

Beyond Linear Algebra

Since GPUs obviously go beyond linear algebra, Thrust and CUB are perfect for exploring other domains in heterogeneous computing. I’ve added snippets that mostly revolve around sorting algorithms, showcasing the differences in memory management between Thrust and CUB and explaining why CUB calls often come in pairs, like:

size_t temp_size = 0;
void *d_temp = nullptr;
cub::DeviceRadixSort::SortKeys(nullptr, temp_size, d_in_keys, d_out_keys, count);
cudaMalloc(&d_temp, temp_size);
cub::DeviceRadixSort::SortKeys(d_temp, temp_size, d_in_keys, d_out_keys, count);

This was also a good place to show how Thrust and CUB operations can be scheduled together on the same asynchronous streams and profiled with GPU time instead of CPU time to avoid unnecessary blocking ⏲️

---

Enjoy exploring, and happy GPU hacking! I’ll keep adding to this project (and other related ones) as we go along!

Changelog

• Add: Binary BMMA kernels for GPU (6a609a0)
• Add: Tensor Core intrinsic benchmarks (1bdb5df)
• Add: cuBLAS benchmarks (2f791fe)
• Add: Precompiled CUDA C++ kernels (c1a6f3e)
• Add: Using CUDA Driver API to JIT .ptx (82cb684)
• Add: PTX and .cuh kernels (824e473)
• Add: Sorting with thrust and cub (df3b2c1)
• Add: Thrust, CUB, CUDA sorting (551402d)
• Add: Thrust, CUB, CUDA sorting (8481114)
• Make: Drop OpenBLAS (3c92c36)
• Fix: Use f16 MMA (141d285)
• Fix: Lower PTX version for JIT (eff3854)
• Fix: Working PTX kernel (514db0f)
• Docs: Introduce Warp-Group-MMA on Hopper (400f294)
• Make: Build CUDA for multiple platforms (3283ab0)
• Fix: Avoid optimizing-out SASS code (986b8bc)
• Fix: Compiling cuBLAS calls (312409a)
• Make: Don't compile PTX (53202e6)
• Make: Silence NVCC warnings (a6cdc74)
• Fix: NVCC compilation issues (494e705)
• Make: Upgrade fmt for NVCC builds (88277bf)
• Fix: Ranges require constexpr on NVCC (c1d7b2f)
• Make: Switch to CUDA Toolkit for GPU libs (2589a40)
• Make: Options for CUDA & TBB in CMake (4d03c08)

v0.5.4: Supporting MSVC on Windows 🪟

The less_slow.cpp project now supports Microsoft Visual C++ (MSVC), thanks to the extensive list of patches suggested by @RazielXYZ 👏

Key updates include switching to OpenBLAS via FetchContent for comparable linear algebra performance across platforms, enabling OpenMP on MSVC for parallelism in Eigen-based computations, and revising OpenMP loop indices to use int64_t, as MSVC requires signed types for parallel loops. The detection of physical cores on high-core-count Windows systems has also been improved by implementing GetActiveProcessorCount(ALL_PROCESSOR_GROUPS) and refining physical core detection logic. Furthermore, the integration addresses missing functionality, such as the lack of __builtin_popcountll on MSVC, with a manual fallback for is_power_of_two.

Additional findings include MSVC-specific behaviors, such as how linking AVX-512 code significantly slows down builds and how assembly-based benchmarks require further investigation for proper MSVC integration. Interestingly, heavily-templated libraries, like Ranges-v3 or CRTE (Compile-time RegEx), show much worse performance on MSVC than with GCC and Clang.

Release v0.5.3

Release: v0.5.3 [skip ci]

Patch

• Improve: Sorting includes (dfbfaa2)
• Make: Recommend VS Code extensions (776b7a4)
• Fix: edge_t constructor for NVCC (5b5c464)

Release v0.5.2

Release: v0.5.2 [skip ci]

Patch

• Improve: Mark non-executable stack (65bac92)

Release v0.5.1

Release: v0.5.1 [skip ci]

Patch

• Improve: MinWarmUpTime before spread ops (f3473c3)
• Docs: Listing Google Benchmark APIs (a19fef3)
• Improve: Use aligned_array in more places (1539624)
• Improve: _k suffix for constants (53f9e71)

Release v0.5.0

Release: v0.5.0 [skip ci]

Minor

• Add: Unbundled FMA versions (5d206f1)

Patch

• Improve: Reorder MA kernels (c3f3ce9)
• Fix: Handling denormal values in FMA (f9b989a)
• Docs: Graviton 4 GEMM benchmarks (108ae72)
• Docs: Note on __bf16 on Arm (651d7e6)
• Fix: Call tops_f64_neon_asm_kernel (44f8086)

Release v0.4.2

Release: v0.4.2 [skip ci]

Patch

• Docs: Better highlights (c662a6c)

v0.4: Allocators and Sparse Graphs 🕸️

This release explores high-performance graph implementations optimized for different access patterns, focusing on real-world use cases like recommendation systems and social networks. To implement those, various STL and Abseil-based containers are used to implement sparse Graph structures.
It shows:

• How to use STL's polymorphic allocators?
• How do we construct a hybrid nested container that propagates stateful allocators to the inner structures?
• Up to 300x performance difference between good implementations in different workload patterns!

Of other neat tricks, shows:

• How can the three-way comparison operator be used with std::tie?
• What's the difference between std::weak_ordering and the strong one ?
• Where can the [[no_unique_address]] attribute be used?

Implementation

It extends the Graph API to:

• upsert_edge(from, to, weight): Inserts or updates an existing edge between two vertices.
• get_edge(from, to): Retrieves the std::optional weight of the edge between two vertices.
• remove_edge(from, to): If present, remove the edge between two vertices.
• for_edges(from, visitor): Applies a callback to all edges starting from a vertex.
• size(): Returns the graph's number of vertices and edges.
• reserve(capacity): Reserves memory for the given number of vertices.
• compact(): Compacts the memory layout of the graph, preparing for read-intensive workloads.

Results

On Intel Sapphire Rapids CPUs in AWS c7i instances:

------------------------------------------------------------------------------------------
Benchmark                                                Time             CPU   Iterations
------------------------------------------------------------------------------------------
graph_make<std::unordered_maps>/min_time:10.000   57329425 ns     57327619 ns          245
graph_make<std::map>/min_time:10.000             109704078 ns    109697310 ns          100
graph_make<absl::flat_set>/min_time:10.000        80598043 ns     80595813 ns          174
graph_rank<std::unordered_maps>/min_time:10.000   35763632 ns     35762406 ns          392
graph_rank<std::map>/min_time:10.000              51658552 ns     51657290 ns          271
graph_rank<absl::flat_set>/min_time:10.000          236938 ns       236933 ns        59137

On AWS Graviton 4 CPUs in AWS r8g instances:

------------------------------------------------------------------------------------------
Benchmark                                                Time             CPU   Iterations
------------------------------------------------------------------------------------------
graph_make<std::unordered_maps>/min_time:10.000  163664945 ns    163660572 ns           86
graph_make<std::map>/min_time:10.000             382543113 ns    382534380 ns           45
graph_make<absl::flat_set>/min_time:10.000       213277341 ns    213272284 ns           64
graph_rank<std::unordered_maps>/min_time:10.000   59579530 ns     59578435 ns          240
graph_rank<std::map>/min_time:10.000              69177429 ns     69175965 ns          191
graph_rank<absl::flat_set>/min_time:10.000          186428 ns       186430 ns        74929

v0.4: Pointer tagging 🏷️

This release provides a prototype of a memory-tagging allocator arena, documenting the complexity of using memory tagging techniques even on Linux:

• Intel's Linear Address Masking
• AMD's Upper Address Ignore
• ARM's Top Byte Ignore
• ARM's Memory Tagging Extension

Minor

• Add: Intel's Linear Address Masking (ead0cd2)
• Add: Pointer tagging draft (72dfc31)

Patch

• Docs: Arm & AMD pointer tagging (3a4b1d7)
• Fix: Avoid tagging if Arm or LA57 (5ac64d1)
• Improve: Log mean allocation size (83ba346)