Quickstart#

A minimal end-to-end example: include the umbrella header, write a kernel that calls a GLASS function, and launch one block per data item.

The kernel#

#include "glass.cuh"

__global__ void my_kernel(float* A, float* B, float* C, int m, int n, int k) {
    // Runtime size: all threads in the block cooperate
    glass::gemm(m, n, k, 1.f, A, B, 0.f, C);
}

Launch with one block per data item:

my_kernel<<<num_items, 256>>>(A, B, C, m, n, k);

That’s the whole contract: every GLASS function assumes it runs inside one CUDA block, and you launch one block per independent problem.

Compiling#

The pure-SIMT path is header-only — just add the repository root to your include path:

nvcc -std=c++17 -I /path/to/GLASS -arch=sm_86 my_kernel.cu -o my_kernel

Compile-time sizes#

Passing the sizes as template arguments lets the compiler unroll the inner loops — the best choice for small fixed-size matrices:

#include "glass.cuh"

__global__ void k(float* A, float* B, float* C) {
    // Sizes baked in as template params — compiler can unroll loops
    glass::gemm<float, 6, 6, 6>(1.f, A, B, 0.f, C);
    glass::gemv<float, 6, 6>(1.f, A, B, 0.f, C);
    glass::axpy<float, 36>(1.5f, A, B);
}

A few more vector/matrix calls (runtime sizes shown):

glass::gemm(m, n, k, 1.f, A, B, 0.f, C);   // C = alpha*A*B + beta*C
glass::gemv(m, n, 1.f, A, x, 0.f, y);       // y = alpha*A*x + beta*y
glass::axpy(n, 1.5f, x, y);                 // y = alpha*x + y

Note

Matrices default to column-major (Fortran) order, consistent with cuBLAS. gemm follows the standard BLAS convention — glass::gemm<T, M, N, K, TRANSPOSE_A, TRANSPOSE_B, ROW_MAJOR_C> (C is M×N, contraction K) — and has no per-operand row-major flag: a row-major M×K operand occupies the same bytes as a column-major K×M matrix, so you just read it with TRANSPOSE_*=true (the single ROW_MAJOR_C flag covers the output). gemv is the exception and keeps a per-matrix ROW_MAJOR flag, because its TRANSPOSE selects the mathematical op (A·x vs Aᵀ·x) and cannot also stand in for storage order. See the “row-major is just a transpose” walkthrough and convention primer in examples/02_gemm_conventions.cu.

Next steps#

  • Using the NVIDIA Backend — route to cuBLASDx / cuSOLVERDx for larger shapes.

  • Library Overview — the call surfaces (block-scoped backends + warp-scoped) and when to use each.

  • Concepts — backend dispatch, TRAILING_SYNC, tuning, and batched-1D APIs.