---
title: Bandicoot GPU Linear Algebra Toolkit
url: https://www.emergentmind.com/topics/bandicoot-gpu-linear-algebra-toolkit
type: topic
---

# Bandicoot GPU Linear Algebra Toolkit

Bandicoot is a C++ linear algebra toolkit that enables high-performance GPU-accelerated computation using a template-based compile-time expression fusion strategy. The toolkit is designed to offer a user-friendly, Armadillo-compatible API, require minimal code modifications for GPU migration, and rival or exceed the efficiency of vendor-optimized and JIT-based frameworks. Bandicoot provides dense and hybrid sparse support, automatic expression kernel fusion at compile-time, backend modularity (CUDA, OpenCL, and planned HIP/Vulkan), and significant empirical advantages over mainstream CPU and GPU alternatives [2604.22242], [2508.11385], [2308.03120], [1904.03383].

## 1. Design Principles and API Compatibility

Bandicoot’s central goal is high efficiency without sacrificing ergonomic usability. It achieves complete API compatibility with Armadillo, the popular CPU linear algebra library: types such as `coot::fmat`, `coot::vec`, and methods like `trans()`, `accu()`, and operator overloading mirror the Armadillo interface. Transitioning an Armadillo-based codebase to Bandicoot generally requires only substitutions of include paths and namespaces:

```cpp
#include <armadillo>  // becomes  #include <coot>
using namespace arma;// becomes  using namespace coot;
```

Matrices and vectors created in this manner reside directly in GPU memory, and the API preserves both the operator syntax and delayed-evaluation semantics of Armadillo. This allows direct reuse of mathematical syntax and high-level algorithms in classic domains ranging from scientific computing to large-scale machine learning [2604.22242], [2508.11385], [2308.03120].

## 2. Software Architecture and Internal Mechanisms

Bandicoot employs a layered architecture:

1. **User API Layer:** Armadillo-compatible matrix, vector, and sparse types (`coot::fmat`, `coot::vec`, `coot::sp_fmat`).
2. **Expression-Building Layer:** Compound expressions are constructed as C++ template ASTs using internal types (e.g., `Op`, `eOp`, `Glue`, `eGlue`), capturing operations such as transpose, multiplication, element-wise transformation, and fusing them before kernel selection.
3. **Kernel-Generation Layer:** Specialized skeleton kernels (e.g., for copy, element-wise apply, GEMM) are macro-instantiated for each AST at compile time. Each compound expression is fused into a unique GPU kernel at compile-time, avoiding runtime code generation or JIT [2604.22242].
4. **Backend Layer:** Modular support for device APIs, currently mature for CUDA and OpenCL, with planned support for HIP, Vulkan, and Metal; these manage device selection, memory allocation, kernel launching, and synchronization [2508.11385].
5. **Utility Layer:** Meta-programming traits, helpers, and compile-time optimization patterns.

This architecture ensures aggressive elimination of intermediate temporaries, and efficient memory bandwidth utilization.

## 3. Compile-Time Expression Fusion and Kernel Generation

Every expression, such as $Y = X.t() + 3$ or $Z = A + B * C - D$, is encoded as nested C++ template types (`eOp`, `Glue`, etc.), forming a full AST at compile time. A recursive metaprogramming traversal extracts the input operands, generates specialized macros for argument access (e.g., `COOT_OBJECT_i_AT(row,col)`), and emits a skeleton kernel body with fully inlined operations:

```cpp
__global__ void coot_kernel_Expr1(float* Z, const float* A, const float* B, const float* C, const float* D, size_t n_rows, size_t n_cols, …) {
  size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < n_rows * n_cols) {
    row = idx % n_rows;
    col = idx / n_rows;
    float vB = B[row + col*B_rows];
    float vC = C[row + col*C_rows];
    float vD = D[row + col*D_rows];
    Z[idx] = A[idx] + (vB * vC) - vD;
  }
}
```

Kernels launch with a default 1-D grid/block partitioning (256 threads per block), using coalesced, strided memory access ($idx = row + col \cdot n_\mathrm{rows}$) to maximize bandwidth on CUDA hardware. The full compound expression—regardless of depth or mixing of ops (add, multiply, transpose, element-wise functions)—is implemented in a single pass, minimizing device memory traversal [2604.22242].

## 4. GPU Integration, Backend Modularity, and Memory Management

Bandicoot supports both CUDA and OpenCL as first-class backends, with device selection handled explicitly or automatically. Core BLAS and LAPACK functionality (GEMM, LU, SVD, eigensolvers, etc.) is mapped onto cuBLAS/cuSolver/cuRand for CUDA or clBLAS/clMAGMA for OpenCL, ensuring broad compatibility.

Device memory allocation is contiguous, and all matrix/vector operations are performed exclusively on the GPU. Host-device conversions require explicit calls, e.g.,

```cpp
arma::fmat H = conv_to<arma::fmat>::from(X_gpu); // GPU→CPU
X_gpu = conv_to<coot::fmat>::from(H);           // CPU→GPU
```

Accessing individual elements on the host triggers an expensive transfer; thus, design discourages per-element reads/writes and forbids iterators. Synchronization is explicit (e.g., `coot_synchronise()`) but kernel launches are asynchronous with respect to the host, allowing pipelining within each operation [2508.11385], [2308.03120].

## 5. Performance, Empirical Evaluation, and Benchmarks

On modern hardware such as NVIDIA RTX 4090, Bandicoot attains effective memory-bandwidth up to 820 GB/s for wide-bandwidth operations, routinely reaching >92% of the device’s practical throughput (868 GB/s measured via `nvbandwidth`). Key empirical findings:

- In multi-matrix addition chains ($Z = M_1 + M_2 + \cdots + M_k$, $k$ up to 16), Bandicoot maintains peak throughput, where other toolkits’ performance degrades sharply for $k>8$ due to increased kernel launches [2604.22242].
- Across 14 tested expressions (addition, chain multiplication, submatrix ops, element-wise fusions, activations, etc.), Bandicoot is fastest or within 5% of the fastest toolkit on all but one benchmark. There is no runtime JIT overhead; all kernels are program-compiled.
- On large practical problem sizes ($>10^3$ per dimension), Bandicoot routinely delivers 10–100× speedups over OpenBLAS on CPU. Small problem sizes (<$10^5$ elements) may run slower due to kernel launch/data movement overhead.
- For GEMM and strided GEMM (where other libraries provided no optimized kernel), Bandicoot delivers up to 66× speedup over naïve fallbacks [1904.03383].

## 6. Mathematical Underpinnings and Optimization Strategy

Bandicoot’s compile-time fusion, enabled by C++ template metaprogramming and partial specialization (SFINAE), covers the following:

- Expressions such as $C = A \times B$ are detected and mapped to BLAS routines or fused kernels with optimal flags (e.g., GEMM with transpose).
- Redundant operations are pruned at compile time (e.g., $\mathrm{trans}(\mathrm{diagmat}(V)) \to \mathrm{diagmat}(V)$).
- Multi-type and element-wise operations are collapsed into minimal, high-throughput kernels.
- Performance is modeled using a bandwidth roofline: for memory-bound ops, $B_\mathrm{saturate} = \frac{R + W}{t_\text{runtime}}$ ($R$: bytes read, $W$: bytes written). This analytic model informs auto-tuned decisions in the Bandicoot search and code generation process [2604.22242], [1904.03383].

For backend code generation, Bandicoot can explore a multi-dimensional optimization space (loop ordering, tiling size, vectorization, memory placements, mapping of logical tiles to threads/blocks) as a constraint satisfaction problem (CSP), using Monte Carlo Tree Search and branch-and-bound strategies [1904.03383].

## 7. Limitations, Extensions, and Future Trajectory

Current constraints include support only for 32-bit float matrices, mature backends limited to CUDA and OpenCL, partial API coverage for advanced factorizations (SVD, Cholesky, eigen), and non-trivial latency for random-access element read/write. Planned extensions are multi-type and low-precision support (FP16, BF16, fixed-point), porting of full LAPACK-style functionality, mixed-precision solvers with GPU tensor core utilization, multi-GPU support, and enhanced auto-tuning of kernel block sizes [2604.22242].

## 8. Installation, Licensing, and Ecosystem Integration

Bandicoot requires a C++11 (or later) compiler, CUDA Toolkit, or OpenCL headers and appropriate BLAS/LAPACK libraries. Linking is simplified via a single plugin library (`-lbandicoot`), abstracting away backend dependencies. The codebase is open-source (Apache 2.0 license) and available at:

```
https://coot.sourceforge.io
```

Bandicoot’s API design promotes easy integration into algorithmic libraries for numerical optimization and machine learning, with ongoing efforts to support end-to-end GPU acceleration for projects such as ensmallen and mlpack [2508.11385], [2308.03120].

---

Bandicoot exemplifies a template-metaprogramming paradigm in C++ GPU linear algebra, where compile-time construction of expression ASTs and macro-driven kernel fusion yield single-pass, memory-bandwidth-saturating computation with no runtime or JIT overhead. Its Armadillo-compatible interface, modular backend support, and open-source licensing position it as an efficient and practical platform for modern large-scale scientific and engineering workflows [2604.22242], [2508.11385], [2308.03120], [1904.03383].

Source: https://www.emergentmind.com/topics/bandicoot-gpu-linear-algebra-toolkit