Skip to content

nvidia.rst

Latest commit

 

History

History
84 lines (60 loc) · 2.33 KB

nvidia.rst

File metadata and controls

84 lines (60 loc) · 2.33 KB
Raw

NVIDIA Backends (glass::nvidia::*)

Vendor-accelerated paths built on CUB (reductions), cuBLASDx (GEMM/GEMV), and cuSOLVERDx (LAPACK). Block-scope ops live in glass::nvidia::block::. Those entry points auto-dispatch at compile time between a pure-SIMT implementation and the vendor backend based on a size heuristic / tuning table (a constexpr decision — no runtime branching); see :doc:`../user_guide/concepts/backend_dispatch`. The L2/L3/LAPACK paths require NVIDIA MathDx (MATHDX_ROOT) — see :doc:`../user_guide/getting_started/installation`.

cuSOLVERDx 0.4+ also supplies explicit smem-less LAPACK wrappers under glass::nvidia::thread::. Each calling CUDA thread owns one packed problem. Warp-scope CUB reductions live in glass::nvidia::warp::. There is no bare operation re-export: the scope is always explicit because it determines launch geometry.

Each call has a companion host-side query helper (*_scratch_bytes, *_threads, *_block_threads_valid) used to size the launch.

L1 (CUB-backed reductions)

reduce / dot / nrm2 at block scope (cub::BlockReduce) and, in glass::nvidia::warp::, at warp scope (cub::WarpReduce — one FULL 32-lane warp per problem, per-warp scratch via warp_reduce_scratch_bytes<T>(), TRAILING_SYNC emits __syncwarp()).

.. doxygenfile:: src/nvidia/l1.cuh
   :no-link:

L2 (gemv)

.. doxygenfile:: src/nvidia/l2.cuh
   :no-link:

L3 (gemm)

.. doxygenfile:: src/nvidia/l3.cuh
   :no-link:

L3 SIMT batched (no cuBLASDx)

.. doxygenfile:: src/nvidia/l3_simt.cuh
   :no-link:

LAPACK (cuSOLVERDx)

.. doxygenfile:: src/nvidia/lapack.cuh
   :no-link:

LAPACK thread execution (cuSOLVERDx 0.4+)

.. doxygenfile:: src/nvidia/lapack_thread.cuh
   :no-link:

Dispatch & query helpers

.. doxygenfile:: src/nvidia/query.cuh
   :no-link:


.. doxygenfile:: src/nvidia/query_simt.cuh
   :no-link:


.. doxygenfile:: src/nvidia/types.cuh
   :no-link:

Host helpers

The umbrella headers also expose host-callable helpers for sizing dynamic shared memory at launch time.

.. doxygenfunction:: glass_gemm_dispatch_smem


.. doxygenfile:: glass-nvidia.cuh
   :no-link: