NumPy-like tensor expressions. C++ control. NVIDIA GPU speed.
Get started · Read the docs · Explore examples · Join the discussion
MatX is a C++20 library for general numerical computing on NVIDIA GPUs and CPUs. Think NumPy-style array programming with native C++ control and speed: create arbitrary-rank tensors, compose lazy expressions, and choose an optimized GPU or multithreaded CPU executor when the work runs.
- 🧮 Write familiar array expressions — broadcasting, slicing, reductions, transforms, linear algebra, solvers, random numbers, and more.
- ⚡ Fuse compatible numerical pipelines — eliminate launches and intermediate memory traffic by changing the executor, not the algorithm.
- 🚀 Reach NVIDIA GPU speed with less code — use optimized CUDA libraries or generate fused kernels at runtime.
- 🖥️ Run on CPUs too — target multithreaded host execution with NVPL, FFTW, OpenBLAS, or BLIS.
#include <matx.h>
constexpr matx::index_t batch = 1024;
constexpr matx::index_t inputs = 256;
constexpr matx::index_t outputs = 128;
auto x = matx::make_tensor<float>({batch, inputs});
auto weights = matx::make_tensor<float>({inputs, outputs});
auto bias = matx::make_tensor<float>({outputs});
auto activations = matx::make_tensor<float>({batch, outputs});
// bias broadcasts naturally across the batch dimension.
auto layer = matx::tanh(matx::matmul(x, weights) + bias);
(activations = layer).run(matx::cudaExecutor{});This is ordinary numerical array code: a batched matrix product, broadcasting, and an elementwise activation. The same expression model covers slicing, reshaping, generators, reductions, convolutions, sparse tensors, linear algebra, solvers, random numbers, and FFTs. Change the executor to target CUDA, JIT fusion, or optimized CPU execution where supported.
| Familiar NumPy idea | MatX model |
|---|---|
ndarray and views |
Arbitrary-rank tensors, non-owning views, slicing, reshaping, and permutation |
| Array creation | zeros, ones, eye, linspace, random generators, and custom generators |
| Vectorized expressions | Lazy elementwise operators, automatic broadcasting, and batched transforms |
| Reductions and numerical routines | Reductions, BLAS, sparse operations, FFTs, decompositions, and solvers |
| CPU/GPU selection | Run the same expression with a CUDA, JIT, or host executor |
FFT processing is a useful fusion example because it crosses an optimized-library boundary and makes the cost of intermediate tensors easy to measure:
#include <matx.h>
constexpr matx::index_t batches = 256;
constexpr matx::index_t samples = 4096;
auto signal = matx::make_tensor<matx::fcomplex>({batches, samples});
auto calibration = matx::make_tensor<float>({batches, samples});
auto spectrum_db = matx::make_tensor<float>({batches, samples});
auto normalized_signal = signal / (matx::abs(signal) + 1.0e-6f);
auto pipeline = 20.0f * matx::log10(
matx::abs(matx::fft(normalized_signal) * calibration) + 1.0e-6f);
(spectrum_db = pipeline).run(matx::CUDAJITExecutor{});The syntax reads like the algorithm: normalize every complex input sample, transform every batch, apply frequency-domain calibration, compute magnitude, and convert the result to decibels. With -DMATX_EN_MATHDX=ON and a supported FFT shape, MatX lowers the compatible pipeline into one generated kernel—no normalized-input or FFT temporary, and no separate pre- or post-processing launch. The same fusion model also extends to supported matrix multiplication, solver, and random-number expressions.
For memory-bound workloads, eliminating launches and intermediate reads and writes lets a simple expression reach speed-of-light performance without hand-writing CUDA kernel plumbing.
Measured on an NVIDIA DGX Spark: 5 lines of MatX replace 59 lines of CUDA + cuFFT and ran 1.61× faster with JIT enabled.
Actual performance depends on the operation, shape, data type, and target GPU.
CUDAJITExecutor can fuse supported transforms and their surrounding operators into a single runtime-generated kernel. With -DMATX_EN_MATHDX=ON, compatible paths include:
| What can join the expression | |
|---|---|
| cuBLASDx | Compatible matrix multiplication and GEMM paths |
| cuSolverDx | Cholesky, inverse, and selected LU, QR, and eigensolver projections |
| cuFFTDx | 1D FFTs and supported single-block 2D complex FFTs |
| cuRANDDx | Small floating-point and complex random expressions |
JIT kernel fusion is experimental and intentionally rejects unsupported combinations rather than silently changing the execution model. See the fusion guide for the current support matrix, constraints, and compile-cache behavior.
The reproducible FFT benchmark example compares handwritten CUDA + cuFFT with MatX cudaExecutor. Enable MathDx to add the JIT case:
- Handwritten CUDA normalization and post-processing kernels around a batched cuFFT plan.
- One MatX expression with
cudaExecutor, which uses cuFFT plus generated normalization and post-processing kernels. - The exact same MatX expression with
CUDAJITExecutor, which fuses both sides of the FFT through cuFFTDx.
- ✂️ 92% less implementation code: 5 MatX lines replace 59 lines of CUDA + cuFFT.1
- ⚖️ Near handwritten CUDA speed: the concise
cudaExecutorpath lands within about 3% of the handwritten implementation. - ⚡ Faster with JIT fusion: the same five-line expression runs 1.61× faster by changing only the executor.2
| Implementation | Median time | Throughput | vs. CUDA + cuFFT |
|---|---|---|---|
| CUDA kernels + cuFFT | 0.1792 ms | 5.851 Gsample/s | 1.00× |
MatX cudaExecutor |
0.1848 ms | 5.675 Gsample/s | 0.97× |
MatX CUDAJITExecutor |
0.1114 ms | 9.413 Gsample/s | 1.61× |
The JIT path removes roughly 68 µs—or 38%—from every pipeline execution.
See the implementations and reproduce the benchmark
The timed handwritten path requires two custom kernels plus explicit FFT planning, two temporary tensors, launch geometry, stream wiring, and error handling:
CUDA + cuFFT
__global__ void NormalizeKernel(/* ... */) { /* magnitude + complex divide */ }
__global__ void MagnitudeDbKernel(/* ... */) { /* calibrate + magnitude + dB */ }
cufftComplex *normalized;
cufftComplex *fft_output;
cudaMalloc(&normalized, count * sizeof(cufftComplex));
cudaMalloc(&fft_output, count * sizeof(cufftComplex));
cufftPlanMany(&plan, 1, size, /* layout... */, CUFFT_C2C, batches);
cufftSetStream(plan, stream);
NormalizeKernel<<<blocks, threads, 0, stream>>>(normalized, signal, count);
cufftExecC2C(plan, normalized, fft_output, CUFFT_FORWARD);
MagnitudeDbKernel<<<blocks, threads, 0, stream>>>(
output, fft_output, calibration, count);The two MatX versions share the same readable expression. Only the executor changes:
MatX CUDA and JIT
auto normalized_signal = signal / (matx::abs(signal) + 1.0e-6f);
auto pipeline = 20.0f * matx::log10(
matx::abs(matx::fft(normalized_signal) * calibration) + 1.0e-6f);
// cuFFT performance with far less code
(spectrum_db = pipeline).run(matx::cudaExecutor{stream});
// Same expression, now fused into one kernel
(spectrum_db = pipeline).run(matx::CUDAJITExecutor{stream});The table reports the median of five benchmark executions. Each execution uses one stream, cached plans, 10 warmups, 2,000 iterations per trial, and seven interleaved trials. Both MatX paths stayed within the benchmark's relative-error tolerance. Online JIT compilation is measured separately and excluded from steady-state timing because its cost depends heavily on the local compile cache.
To reproduce the comparison on your GPU:
cmake -S . -B build \
-DCMAKE_BUILD_TYPE=Release \
-DMATX_BUILD_EXAMPLES=ON \
-DMATX_EN_MATHDX=ON
cmake --build build --target fft_benchmark -j
./build/examples/fft_benchmark 256 4096 2000fft_benchmark is built with the other examples whenever MATX_BUILD_EXAMPLES=ON. The JIT case is compiled only with MATX_EN_MATHDX=ON. If MatX is configured with MATX_EN_NVPL=ON or MATX_EN_X86_FFTW=ON, the example also compiles and reports the identical expression through the multithreaded AllThreadsHostExecutor. That optional CPU result is printed separately from the GPU cases.
Performance depends on GPU, clocks, shape, and software versions; use the benchmark rather than assuming this ratio for every workload.
MatX is not GPU-only. The same operator model runs on CPUs, and AllThreadsHostExecutor uses OpenMP to distribute supported work across the available host threads:
matx::AllThreadsHostExecutor host{};
(activations = layer).run(host);Host transforms dispatch to optimized CPU libraries when their corresponding CMake option is enabled:
| Backend | Optimized host operations | Enable with |
|---|---|---|
| NVIDIA NVPL | FFT, BLAS, and LAPACK on NVIDIA ARM CPUs | -DMATX_EN_NVPL=ON |
| FFTW | FFT on x86 CPUs | -DMATX_EN_X86_FFTW=ON |
| OpenBLAS | BLAS and LAPACK on x86 CPUs | -DMATX_EN_OPENBLAS=ON |
| BLIS | BLAS on x86 CPUs | -DMATX_EN_BLIS=ON |
All elementwise operators and reductions have host paths, as do FFT, BLAS, and LAPACK transforms when the relevant backend is configured. Most host functions support multithreading; reductions and a small set of transforms currently remain single-threaded. See executor compatibility for the operation-by-operation breakdown.
The Black–Scholes example expresses the complete pricing equation in concise MatX syntax and evaluates it as one fused GPU kernel—without hand-written CUDA or materialized intermediate arrays.
For 100 million independent FP32 option prices on the same NVIDIA DGX Spark, the fused MatX expression is 5.62× faster than the equivalent eager CuPy expression on the same GPU and 184.5× faster than NumPy/SciPy.3
| Implementation | Median time | Throughput | vs. NumPy |
|---|---|---|---|
| NumPy 2.5.1 + SciPy 1.18.0 / CPU | 2965.66 ms | 33.7 M options/s | 1.0× |
| CuPy 14.1.1 eager expression / GPU | 90.32 ms | 1.107 B options/s | 32.8× |
| MatX expression / GPU | 16.07 ms | 6.223 B options/s | 184.5× |
Built for real time. At 6.223 billion results per second, even a single FP32 output stream represents nearly 200 Gbps—giving MatX the speed to process data in real time at hundreds of gigabits per second.
Input creation and transfers are excluded. NumPy and CuPy report the median of five trials after two warmups. The MatX result is the median of five runs of the existing example; each run averages 100 CUDA-event-timed MatX iterations. The companion NumPy/CuPy benchmark keeps the equation, FP32 dtype, array size, and timing boundaries reproducible.
Reproduce the Black–Scholes comparison
cmake -S . -B build -DMATX_BUILD_EXAMPLES=ON
cmake --build build --target black_scholes -j
./build/examples/black_scholes
python -m pip install numpy scipy cupy-cuda13x
python examples/black_scholes_benchmark.py \
--size 100000000 --warmups 2 --trials 5MatX is header-only. Add the repository to your project and link its interface target:
cmake_minimum_required(VERSION 3.30.4)
project(my_app LANGUAGES CXX CUDA)
add_subdirectory(path/to/MatX)
add_executable(my_app main.cu)
target_link_libraries(my_app PRIVATE matx::matx)Then include the API:
#include <matx.h>To build MatX examples and tests locally:
cmake -S . -B build \
-DMATX_BUILD_EXAMPLES=ON \
-DMATX_BUILD_TESTS=ON
cmake --build build -j
ctest --test-dir buildMatX currently targets Linux and requires a C++20 build environment. GPU builds require the CUDA Toolkit and a supported host compiler; consult the build guide for exact current versions, optional backends, offline setup, and CPU support.
| I want to… | Go here |
|---|---|
| Learn the tensor and expression model | Quick start |
| Understand kernel fusion | Fusion guide |
| Translate familiar array syntax | MATLAB / Python syntax guide |
| Check executor support | Executor compatibility |
| Learn interactively | Self-guided notebooks |
| Start from working applications | Examples |
- Ask usage questions in GitHub Discussions.
- Report bugs and request features through GitHub Issues.
- Read CONTRIBUTING.md before submitting a change.
- MatX uses semantic versioning and is available under the BSD 3-Clause license.
Footnotes
-
Counted from the benchmark source, excluding comments, blank lines, shared tensor/input setup, timing, and validation. The CUDA + cuFFT count includes both custom kernels, temporary allocation and cleanup, FFT planning, launch geometry, and launches; the MatX count includes the expression, executor, and execution. ↩
-
Measured with 256 batched 4096-point complex FP32 FFTs on an NVIDIA DGX Spark (GB10) using CUDA 13.0 and MathDx 26.03. ↩
-
Measured on an NVIDIA DGX Spark (GB10) with CUDA 13.0 and driver 580.159.03, including its 20-core Arm CPU (10 Cortex-X925 + 10 Cortex-A725 cores). ↩