Using the NVIDIA Backend#

The glass::nvidia::block:: interface routes to NVIDIA’s device-side libraries — CUB (L1), cuBLASDx (L2/L3 GEMM/GEMV/batched), and cuSOLVERDx (LAPACK) — while preserving the same one-block __device__ calling convention. These wrappers require compile-time matrix sizes. (Bare glass::nvidia:: spellings are the measured-default face and currently resolve to the same entities — see Namespaces, suffixes, and flags. There is also glass::nvidia::warp:: — CUB WarpReduce reduce / dot / nrm2, one FULL 32-lane warp per problem, per-warp scratch via warp_reduce_scratch_bytes<T>(), TRAILING_SYNC emitting __syncwarp().)

Make sure MathDx is installed and MATHDX_ROOT is set first — see Installation.

Compiling against MathDx#

cuBLASDx is header-only; cuSOLVERDx links a precompiled device fatbin. The flags differ by which level you use.

cuBLASDx (GEMM / GEMV / batched):

nvcc -std=c++17 -arch=sm_86 \
     -I /path/to/GLASS \
     -I $MATHDX_ROOT/include \
     -I $MATHDX_ROOT/external/cutlass/include \
     --expt-relaxed-constexpr \
     -Xptxas -O1 \
     -DGLASS_BENCH_CUBLASDX \
     my_kernel.cu -o my_kernel

  • --expt-relaxed-constexpr is required for cuBLASDx headers.

  • -Xptxas -O1 works around the CUDA 12.9 cuBLASDx miscompilation (and doubles as anti-DSE).

  • -DGLASS_BENCH_CUBLASDX force-enables the wrappers when the TU hasn’t pre-included the cuBLASDx header. (GLASS otherwise auto-detects via GLASS_HAVE_CUBLASDX from include order.)

cuSOLVERDx (LAPACK) additionally needs the device-link flags:

nvcc -std=c++17 -arch=sm_86 \
     -I /path/to/GLASS -I $MATHDX_ROOT/include \
     --expt-relaxed-constexpr -Xptxas -O1 \
     -rdc=true -dlto \
     -L $MATHDX_ROOT/lib \
     -lcusolverdx -lcublas -lcusolver -lcudart \
     my_kernel.cu -o my_kernel

SMS defaults to 860 and can be overridden with -DSMS=XXX so the dispatch heuristic and cuBLASDx code-gen target your arch.

Calling glass::nvidia::block:: — default form#

In the default form cuBLASDx picks the thread count; you query the required shared memory and thread count (both constexpr) on the host and launch with the exact thread count — a mismatch deadlocks.

#include "glass-nvidia.cuh"

constexpr auto smem    = glass::nvidia::block::gemm_scratch_bytes<float, 6, 6, 6>();
constexpr auto threads = glass::nvidia::block::gemm_threads<float, 6, 6, 6>();

__global__ void k(float* A, float* B, float* C) {
    extern __shared__ __align__(16) char smem_buf[];
    glass::nvidia::block::gemm<float, 6, 6, 6>(1.f, A, B, 0.f, C, smem_buf);
}

// Launch with the EXACT thread count cuBLASDx wants.
k<<<1, threads, smem>>>(dA, dB, dC);

Caller-pinned BlockDim<TC>#

To launch with a thread count your surrounding kernel needs (e.g. GRiD’s 352-thread launches), pin BlockDim<TC> with a DEFINE macro. Extra threads go idle inside the GEMM:

namespace glass { namespace nvidia { namespace block {
    DEFINE_NVIDIA_GEMM_BLOCKDIM(6, 6, 6, 352)   // pin BlockDim<352,1,1>
}}}

__global__ void k(float* A, float* B, float* C) {
    extern __shared__ __align__(16) char smem_buf[];
    glass::nvidia::block::gemm<float, 6, 6, 6, 352>(1.f, A, B, 0.f, C, smem_buf);
}

constexpr auto smem = glass::nvidia::block::gemm_scratch_bytes<float, 6, 6, 6, 352>();
k<<<1, 352, smem>>>(dA, dB, dC);

Query the smallest valid TC for a (T, M, N, K, SM) tuple:

constexpr uint32_t MIN = glass::nvidia::block::gemm_min_block_threads<float, 6, 6, 6>();
static_assert(glass::nvidia::block::gemm_block_threads_valid<float, 6, 6, 6, 352>(),
              "352 threads should be enough for 6x6x6 on this SM");

Tip

Compile without -DNDEBUG and the wrappers assert(blockDim >= GEMM::block_dim) inside every run() — a misconfigured launch fails with a clean assertion instead of a silent deadlock. The asserts compile out under -DNDEBUG.

Layout / transpose#

glass::nvidia::block::gemm accepts layout LA, LB, LC template parameters (mirroring cuBLASDx’s Arrangement<>) so you can express transpose / row-major storage without falling back to SIMT:

namespace glass { namespace nvidia { namespace block {
    // A · Bᵀ  (B row-major); alias for LB=row_major
    DEFINE_NVIDIA_GEMM_BLOCKDIM_TRANSB(6, 6, 6, 352)
    // Or fully explicit (LA=row, LB=col, LC=col):
    DEFINE_NVIDIA_GEMM_BLOCKDIM_LAYOUT(6, 6, 6, 352, 1, 0, 0)
}}}

Layout arguments are integer literals: 0 = col_major, 1 = row_major.

Adding a custom size#

Pre-instantiated square GEMM/GEMV sizes are 4, 6, 8, 12, 14, 24, 64. For any other size, place the DEFINE macro inside namespace glass::nvidia::block in your .cu file (the specializations must live where the primary templates do — the block namespace):

#include "glass-nvidia.cuh"
namespace glass { namespace nvidia { namespace block {
    DEFINE_NVIDIA_GEMM(16, 16, 16)
    DEFINE_NVIDIA_GEMV(16, 16)
    DEFINE_NVIDIA_GEMM_BLOCKDIM(16, 16, 16, 256)
    DEFINE_NVIDIA_POSV_BLOCKDIM(16, 4, 256)   // SPD solve, 4 RHS
}}}

Every wrapper exposes the same macro family — substitute GEMM / GEMV / GEMM_BATCHED / CHOL / TRSM / POSV / POTRS / GETRF / GETRS / GESV / GEQRF / GELS and append _BLOCKDIM / _LAYOUT / _SM / _BLOCKDIM_SM etc. as needed.

Linear solvers (cuSOLVERDx)#

#include "glass-nvidia.cuh"

namespace glass { namespace nvidia { namespace block {
    DEFINE_NVIDIA_POSV_BLOCKDIM(7, 1, 256)   // 7×7 SPD, 1 RHS, BlockDim<256>
}}}

__global__ void k(float* A, float* b) {
    extern __shared__ __align__(16) char smem_buf[];
    // Solves A·x = b in place: A := L (lower Cholesky), b := x
    glass::nvidia::block::posv<float, 7, 1, 256>(A, b, smem_buf);
}

constexpr auto smem = glass::nvidia::block::posv_scratch_bytes<float, 7, 1, 256>();
k<<<1, 256, smem>>>(dA, db);

Available cuSOLVERDx wrappers: potrf, trsm, posv, potrs, getrf_no_pivot, getrs_no_pivot, gesv_no_pivot, geqrf, gels. All follow the same DEFINE_NVIDIA_<NAME> macro pattern and are not pre-instantiated — call the macro per size you need.

See Backend Dispatch for how the auto-dispatch decides between cuBLASDx and SIMT, and Batched-1D GEMM APIs for the 1D-launch batched GEMM APIs.