NVIDIA cuSolverDx

https://docs.nvidia.com/cuda/cusolverdx/
Admin
NVIDIA cuSolverDx is a library that enables selected dense matrix factorization and solve routines...
Tokens:71,379
Snippets:789
Trust Score:10
Update:3 months ago
Show doc for...
Tokens:
Raw
### Launch POSV Kernel and Host Setup

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/introduction.html

Defines the POSV kernel structure and demonstrates the host-side setup, including memory allocation and kernel invocation.

```cpp
template<class POSV, unsigned int BatchesPerBlock = 1,
         class DataType = typename POSV::a_data_type>
__global__ void posv_kernel(DataType* A, const unsigned int lda_gmem,
                            DataType* B, const unsigned int ldb_gmem,
                            typename POSV::status_type* info,
                            const unsigned int batches) {

    using namespace cusolverdx;
    constexpr auto m    = POSV::m_size;
    constexpr auto nrhs = POSV::k_size;
    const auto     one_batch_size_a_gmem = lda_gmem * m;
    const auto     one_batch_size_b_gmem =
          (arrangement_of_v_b<POSV> == arrangement::col_major) ? ldb_gmem * nrhs : m * ldb_gmem;

    constexpr auto lda_smem = POSV::lda;
    constexpr auto ldb_smem = POSV::ldb;
    constexpr auto one_batch_size_a_smem = lda_smem * m;
    constexpr auto one_batch_size_b_smem =
          (arrangement_of_v_b<POSV> == arrangement::col_major) ? ldb_smem * nrhs : m * ldb_smem;

    extern __shared__ __align__(sizeof(DataType)) cusolverdx::byte shared_mem[];

    DataType* As = reinterpret_cast<DataType*>(shared_mem);
    DataType* Bs = As + one_batch_size_a_smem * BatchesPerBlock;

    const auto batch_idx = blockIdx.x * BatchesPerBlock;
    if (batch_idx >= batches)
        return;

    auto Ag = A + size_t(one_batch_size_a_gmem) * batch_idx;
    auto Bg = B + size_t(one_batch_size_b_gmem) * batch_idx;

    // Load data from global memory to registers
    common::io<POSV, BatchesPerBlock>::load(Ag, lda_gmem, As, lda_smem);
    common::io<POSV, BatchesPerBlock>::load_rhs(Bg, ldb_gmem, Bs, ldb_smem);

    POSV().execute(As, lda_smem, Bs, &info[batch_idx]);

    // store
    common::io<POSV, BatchesPerBlock>::store(As, lda_smem, Ag, lda_gmem);
    common::io<POSV, BatchesPerBlock>::store_rhs(Bs, ldb_smem, Bg, ldb_gmem);
}

 int posv_batched() {
    using POSV = decltype(Size<32 /* = M */, 32 /* = N */, 1 /* = K */>()
                      + Precision<double>()
                      + Type<type::complex>()
                      + Function<function::posv>()
                      + FillMode<lower>()
                      + Arrangement<col_major /* A */, row_major /* B */>()
                      + SM<900>()
                      + Block());

    using data_type = typename POSV::a_data_type;

    constexpr auto m    = POSV::m_size;
    constexpr auto n    = POSV::n_size;
    constexpr auto nrhs = POSV::k_size;
    static_assert(m == n, "posv is for Hermitian positive-definite matrix only");
    constexpr auto lda_smem = POSV::lda;
    constexpr auto ldb_smem = POSV::ldb;

    constexpr auto lda = m;
    constexpr auto ldb = (arrangement_of_v_b<POSV> == arrangement::col_major) ? m : nrhs;

    const auto batches          = 2;
    const auto one_batch_size_A = lda * n;
    const auto one_batch_size_B = m * nrhs;

    std::vector<data_type> A(one_batch_size_A * batches);
    // fill the A matrix

    std::vector<data_type> B(one_batch_size_B * batches);
    // fill the B matrix

    std::vector<int> info(batches, 0);
    data_type*       d_A    = nullptr; /* device copy of A */
    data_type*       d_B    = nullptr; /* device copy of B */
    int*             d_info = nullptr; /* error info */

    cudaMalloc(reinterpret_cast<void**>(&d_A), sizeof(data_type) * A.size());
    cudaMalloc(reinterpret_cast<void**>(&d_B), sizeof(data_type) * B.size());
    cudaMalloc(reinterpret_cast<void**>(&d_info), sizeof(int) * batches);

    cudaMemcpy(d_A, A.data(), sizeof(data_type) * A.size(), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B.data(), sizeof(data_type) * B.size(), cudaMemcpyHostToDevice);

    const auto sm_size = POSV::shared_memory_size;

    //Invokes kernel. Batches per block is 1
    posv_kernel<POSV><<<batches, POSV::block_dim, sm_size>>>(d_A, lda, d_B, ldb, d_info);

    cudaDeviceSynchronize();

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_info);
 }
```

--------------------------------

### Cluster Execution Tile Layout Example

Source: https://docs.nvidia.com/cuda/cusolverdx/get_started/functions/potrf.html

Visual representation of a 4x4 global tile layout for upper fill mode.

```text
| A00  A01  A02  A03 |
|  X   A11  A12  A13 |
|  X    X   A22  A23 |
|  X    X    X   A33 |
```

--------------------------------

### Example Usage of is_supported

Source: https://docs.nvidia.com/cuda/cusolverdx/api/other_traits.html

Demonstrates checking solver support across different CUDA architectures for a defined solver configuration.

```cpp
using namespace cusolverdx;

using Solver = decltype(Size<160>() + Function<potrf>() + Type<type::real>() +
                        Block() + Precision<double>());
cusolverdx::is_supported<Solver, 1000>::value; // true
cusolverdx::is_supported<Solver,  900>::value; // true
cusolverdx::is_supported<Solver,  800>::value; // false
```

--------------------------------

### Configure and Link cuSolverDx in CMake

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/installation.html

Standard setup for finding the package, configuring target properties for optimization, and linking the library.

```cmake
# find cuSolverDx
find_package(mathdx REQUIRED COMPONENTS cusolverdx CONFIG)
# enable LTO in your target
set_target_properties(YourProgram
      PROPERTIES
          CUDA_SEPARABLE_COMPILATION ON
          INTERPROCEDURAL_OPTIMIZATION ON)
# link against mathdx::cusolverdx
target_link_libraries(YourProgram mathdx::cusolverdx)
# or against mathdx::cusolverdx_fatbin
target_link_libraries(YourProgram mathdx::cusolverdx_fatbin)
```

--------------------------------

### Shared Memory Slicing Example

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_shared.html

Demonstrates slicing shared memory into a mix of raw pointers and CuTe tensors using the slice function.

```cpp
// Slice shared memory into tensors and raw arrays (pointers)
auto [s_a_ptr, s_cublasdx_a_tensor, s_cubladx_b_tensor, s_array_ptr] =
    cusolverdx::shared_memory::slice<typename POSVDescription::a_data_type, AInputType, BInputType, float>(
        smem,
        alignof(typename POSVDescription::a_data_type), (lda * POSV::m_size),
        cublasdx::alignment_of_v_a<BLAS>, BLAS::suggest_layout_smem_a(),
        cublasdx::alignment_of_v_b<BLAS>, BLAS::suggest_layout_smem_b(),
        alignof(float), 64
    );
```

--------------------------------

### Calculate Shared Memory Usage

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_shared.html

Example demonstrating how to chain add calls to calculate total shared memory for cuSolverDx, cuBLASDx, and custom arrays.

```cpp
# include "cusolverdx_io.hpp"

auto shared_memory_size =
  cusolverdx::make_shared_storage_calculator()
     // cuSolverDx input matrix
     .add(alignof(typename POSVDescription::a_data_type), POSVDescription::shared_memory_size)
     // A and B matrices for cuBLASDx
     .add(cublasdx::alignment_of_v_a<BLAS>, sizeof(AInputType), BLAS::suggest_layout_smem_b())
     .add(cublasdx::alignment_of_v_b<BLAS>, sizeof(BInputType), BLAS::suggest_layout_smem_a())
     // Custom array of 64 elements
     .add(alignof(float), sizeof(float), 64)
     .get();
```

--------------------------------

### Define Solver Function Descriptor

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/execution_ops.html

Example of creating a Solver function descriptor configured for single-block execution.

```cpp
#include <cusolverdx.hpp>

using namespace cusolverdx;
using Solver = decltype(Size<32>()
                    + Precision<double>()
                    + Type<type::real>()
                    + Function<function::potrf>()
                    + FillMode<lower>()
                    + Arrangement<row_major>()
                    + SM<900>()
                    + Block());
```

--------------------------------

### Specify MathDx Package Path

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/installation.html

Use the PATHS option within find_package to point to a specific installation directory.

```cmake
find_package(mathdx REQUIRED COMPONENTS cusolverdx CONFIG PATHS "<your_directory>/nvidia/mathdx/yy.mm/")
```

--------------------------------

### Define Cluster POTRF Solver

Source: https://docs.nvidia.com/cuda/cusolverdx/0.5.0/api/execution_ops.html

Example configuration for a 288x288 double-precision POTRF solver using the Cluster operator on Hopper architecture.

```cpp
#include <cusolverdx_cluster.hpp>

using namespace cusolverdx;
using Solver = decltype(Size<288>()
                    + Precision<double>()
                    + Type<type::real>()
                    + Function<function::potrf>()
                    + FillMode<fill_mode::lower>()
                    + Arrangement<col_major>()
                    + SM<900>()
                    + Cluster()
                    + BlocksPerCluster<4>()
                    + TileSize<96>()
                    + BlockDim<128>());
```

--------------------------------

### Initialize Shared Storage Calculator

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_shared.html

Creates an instance of the shared storage calculator.

```cpp
shared_storage_calculator make_shared_storage_calculator();
```

--------------------------------

### GELQF Device Execution Methods

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/gelqf.html

Device-side function signatures for computing LQ factorization, with variants for fixed and runtime leading dimensions.

```cpp
__device__ void execute(data_type* A, data_type* tau);
// with runtime lda
__device__ void execute(data_type* A, const unsigned int lda, data_type* tau);
```

--------------------------------

### Constructing an Operation Descriptor

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/description_ops.html

Demonstrates how to combine multiple description operators to define a specific linear system solver configuration.

```cpp
#include <cusolverdx.hpp>

using Solver = decltype(cusolverdx::Size<16, 16, 4>()
              + cusolverdx::Precision<double>()
              + cusolverdx::Type<cusolverdx::type::complex>()
              + cusolverdx::Arrangement<cusolverdx::row_major>()
              + cusolverdx::FillMode<cusolverdx::lower>()
              + cusolverdx::Function<cusolverdx::function::posv>()
              + cusolverdx::SM<900>());
```

--------------------------------

### Querying Solver Description Traits

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/description_traits.html

Demonstrates how to define a solver using description operators and query its properties using trait helper functions.

```cpp
#include <cusolverdx.hpp>
using namespace cusolverdx;

using Solver = decltype(Size<64>()
                + Precision<double>()
                + Function<function::potrf>()
                + FillMode<lower>()
                + SM<900>()
                + Block());

if(cusolverdx::is_complete_solver<Solver>::value) {
  std::cout << "Matrix A (M x N): "
    << cusolverdx::size_of<Solver>::m << " x "
    << cusolverdx::size_of<Solver>::n << std::endl;
}
```

--------------------------------

### Compile and link cuSolverDx projects

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/installation.html

Commands for compiling source files with the cuSolverDx library using either the static library or the fatbin library.

```bash
// When using libcusolverdx.a
nvcc -dlto -std=c++17 -arch sm_XY (...) -I<mathdx_include_dir> <your_source_file>.cu -o <your_binary> -lcusolverdx

// When using libcusolverdx.fatbin
nvcc -dlto -std=c++17 -arch sm_XY (...) -I<mathdx_include_dir> <your_source_file>.cu -o <your_binary> libcusolverdx.fatbin
```

--------------------------------

### Apply Suggested Batches Per Block Trait

Source: https://docs.nvidia.com/cuda/cusolverdx/api/other_traits.html

Demonstrates defining an execution descriptor and incorporating the suggested batches per block trait.

```cpp
using Base = decltype(Size<14, 14>() + Precision<double>() + Type<type::complex>() +
              Block() + Function<posv>() + SM<Arch>());

using POSV = decltype(Base() +
                      BatchesPerBlock<Base::suggested_batches_per_block>() +
                      FillMode<fill_mode::lower>() +
                      Arrangement<arrangement::row_major>());
```

--------------------------------

### Using Wrapper Functions for copy_2d

Source: https://docs.nvidia.com/cuda/cusolverdx/0.5.0/api/other_methods.html

Shows the use of convenient wrapper functions provided in the common device_io header for loading and storing matrix data.

```cpp
// Example 2: Using a set of convenient wrapper functions of copy_2d provided in example/common/device_io.hpp
# include "example/common/device_io.hpp"

common::io<Solver, Solver::BatchesPerBlock>::load_a(Ag, lda_g, As, lda_s);
common::io<Solver, Solver::BatchesPerBlock>::load_b(Bg, ldb_g, Bs, ldb_s);

// run the solver functions
Solver::execute(As, lda_s, ...);

// Store results back to global memory
common::io<Solver, Solver::BatchesPerBlock>::store_a(As, lda_s, Ag, lda_g);
common::io<Solver, Solver::BatchesPerBlock>::store_b(Bs, ldb_s, Bg, ldb_g);
```

--------------------------------

### Solver::execute (geqrf / gelqf)

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/methods.html

Executes QR or LQ factorization.

```APIDOC
## Solver::execute (geqrf / gelqf)

### Description
Performs QR (geqrf) or LQ (gelqf) factorization.

### Method
void Solver::execute(typename Solver::a_data_type* A, typename Solver::a_data_type* tau);
void Solver::execute(typename Solver::a_data_type* A, const unsigned int runtime_lda, typename Solver::a_data_type* tau);
```

--------------------------------

### Using Wrapper Functions for IO

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_methods.html

Utilizes convenient wrapper functions from the common IO header to handle matrix loading and storing operations.

```cpp
// Example 2: Using a set of convenient wrapper functions of copy_2d provided in example/common/device_io.hpp
// example/common/device_io.hpp already includes cusolverdx_io.hpp
# include "example/common/device_io.hpp"

common::io<Solver, Solver::BatchesPerBlock>::load_a(Ag, lda_g, As, lda_s);
common::io<Solver, Solver::BatchesPerBlock>::load_b(Bg, ldb_g, Bs, ldb_s);

// run the solver functions
Solver::execute(As, lda_s, ...);

// Store results back to global memory
common::io<Solver, Solver::BatchesPerBlock>::store_a(As, lda_s, Ag, lda_g);
common::io<Solver, Solver::BatchesPerBlock>::store_b(Bs, ldb_s, Bg, ldb_g);
```

--------------------------------

### Solver::suggested_batches_per_block

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_traits.html

Provides the recommended optimal number of batches per block for a Solver description.

```APIDOC
## Solver::suggested_batches_per_block

### Description
Returns the recommended optimal number of batches per block for a Solver description. This is typically used to configure the BatchesPerBlock operator in an execution descriptor.

### Signature
`unsigned int Solver::suggested_batches_per_block`
```

--------------------------------

### Solver::suggested_block_dim

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_traits.html

Provides the recommended number of threads for a Solver description to achieve optimal performance.

```APIDOC
## Solver::suggested_block_dim

### Description
Returns the recommended number of threads for a Solver description, depending on size, data type, GPU architecture, and batches per thread block. This value is used if the BlockDim operator is not explicitly defined by the user.

### Signature
`dim3(X, Y==1, Z==1) Solver::suggested_block_dim`
```

--------------------------------

### execute

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/gels.html

Device function to solve least square problems using QR or LQ factorization.

```APIDOC
## __device__ void execute(data_type* A, data_type* tau, data_type* B)
## __device__ void execute(data_type* A, data_type* tau, data_type* B, const unsigned int ldb)
## __device__ void execute(data_type* A, const unsigned int lda, data_type* tau, data_type* B)
## __device__ void execute(data_type* A, const unsigned int lda, data_type* tau, data_type* B, const unsigned int ldb)

### Description
Solves overdetermined or underdetermined least square problems by overwriting B with the solution X using QR or LQ factorization of A.

### Parameters
- **A** (data_type*) - Required - Batched M×N matrix. Overwritten by QR or LQ factorization.
- **tau** (data_type*) - Required - Array of size min(M,N) per batch, representing Householder vectors.
- **B** (data_type*) - Required - Batched right-hand side matrix. Overwritten by the solution X.
- **lda** (unsigned int) - Optional - Leading dimension of A.
- **ldb** (unsigned int) - Optional - Leading dimension of B.
```

--------------------------------

### execute

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/modified_lu.html

Performs a batched modified LU factorization on an orthonormal matrix Q.

```APIDOC
## __device__ void execute(data_type* Q, data_type* S)

### Description
Computes the modified LU factorization of an orthonormal matrix Q, overwriting Q with the resulting L and U matrices and storing the diagonal sign matrix in S.

### Parameters
- **Q** (data_type*) - Required - Orthonormal matrix of size M×N per batch. Overwritten with L and U on output.
- **S** (data_type*) - Required - 1D array of size min(M,N) per batch. Contains the diagonal sign matrix on output.

---

## __device__ void execute(data_type* Q, const unsigned int ldq, data_type* S)

### Description
Computes the modified LU factorization of an orthonormal matrix Q with a specified leading dimension for Q.

### Parameters
- **Q** (data_type*) - Required - Orthonormal matrix of size M×N per batch. Overwritten with L and U on output.
- **ldq** (const unsigned int) - Required - Leading dimension of matrix Q.
- **S** (data_type*) - Required - 1D array of size min(M,N) per batch. Contains the diagonal sign matrix on output.
```

--------------------------------

### Implement Cholesky factorization and solve in Warp

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/python_bindings.html

Demonstrates a kernel using Warp's tile-based API to perform Cholesky factorization and solve a linear system.

```python
@wp.kernel
def cholesky(
    A: wp.array2d(dtype=wp_type),
    L: wp.array2d(dtype=wp_type),
    X: wp.array1d(dtype=wp_type),
    Y: wp.array1d(dtype=wp_type),
):

    i, j, _ = wp.tid()

    a = wp.tile_load(A, shape=(TILE, TILE))
    l = wp.tile_cholesky(a)
    wp.tile_store(L, l)

    x = wp.tile_load(X, shape=TILE)
    y = wp.tile_cholesky_solve(l, x)
    wp.tile_store(Y, y)
```

--------------------------------

### void execute(const data_type* A, data_type* B)

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/potrs.html

Solves the linear system AX=B using the Cholesky factorization stored in A, overwriting B with the solution X.

```APIDOC
## __device__ void execute(const data_type* A, data_type* B)

### Description
Solves a system of linear equations AX=B where A contains a Cholesky factorization. The matrix B is overwritten with the solution X.

### Parameters
- **A** (const data_type*) - The batched M×M Hermitian matrix containing the Cholesky factorization.
- **B** (data_type*) - The batched right-hand side matrix of size M×K, which is overwritten by the solution X.
```

--------------------------------

### Solver::execute (ungqr / unglq)

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/methods.html

Executes generation of orthogonal matrices.

```APIDOC
## Solver::execute (ungqr / unglq)

### Description
Generates the orthogonal matrix Q.

### Method
void Solver::execute(typename Solver::a_data_type* A, const typename Solver::a_data_type* tau);
void Solver::execute(typename Solver::a_data_type* A, const unsigned int runtime_lda, const typename Solver::a_data_type* tau);
```

--------------------------------

### Cholesky Factorization Kernel with nvmath-python

Source: https://docs.nvidia.com/cuda/cusolverdx/python_bindings.html

Shows a Cholesky factorization kernel implementation using nvmath-python's CholeskySolver and Numba CUDA JIT.

```python
solver = CholeskySolver(
    size=(32, 32),
    precision=np.float64,
    data_type="complex",
    execution="Block",
    leading_dimensions=(33, 33),
    block_dim=(256, 1, 1),
    fill_mode="upper",
)

n_a = solver.a_size()

@cuda.jit
def kernel(a, info):
    smem_a = cuda.shared.array(n_a, dtype=solver.value_type)

    load_to_shared_strided(a, smem_a, solver.a_shape, solver.a_strides())
    cuda.syncthreads()

    solver.factorize(smem_a, info)
    cuda.syncthreads()

    store_from_shared_strided(smem_a, a, solver.a_shape, solver.a_strides())
```

--------------------------------

### Check solver execution configuration

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/description_traits.html

Verifies if a solver descriptor is configured for execution using both description and execution operators.

```cpp
// bool
cusolverdx::is_solver_execution<Solver>::value
cusolverdx::is_solver_execution_v<Solver>
```

--------------------------------

### Configure TileSize

Source: https://docs.nvidia.com/cuda/cusolverdx/0.5.0/api/execution_ops.html

Sets the tile size for the blocked cluster POTRF algorithm, constrained between 1 and N.

```cpp
cusolverdx::TileSize<unsigned int TS>()
```

--------------------------------

### Retrieve Suggested Batches Per Block

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_traits.html

Provides the optimal number of batches per block for a Solver description, often used to configure the execution descriptor.

```cpp
// unsigned int
Solver::suggested_batches_per_block
```

```cpp
using Base = decltype(Size<14, 14>() + Precision<double>() + Type<type::complex>() +
              Block() + Function<posv>() + SM<Arch>());

using POSV = decltype(Base() +
                      BatchesPerBlock<Base::suggested_batches_per_block>() +
                      FillMode<fill_mode::lower>() +
                      Arrangement<arrangement::row_major>());
```

--------------------------------

### Using suggested_tile_size in a Descriptor

Source: https://docs.nvidia.com/cuda/cusolverdx/api/other_traits.html

Apply the suggested tile size to a descriptor configuration when no explicit TileSize is provided.

```cpp
#include <cusolverdx_cluster.hpp>

using namespace cusolverdx;

using Base = decltype(Size<288>() + Precision<float>() + Type<type::real>() +
                      Function<function::potrf>() + FillMode<fill_mode::lower>() +
                      Cluster() + BlocksPerCluster<4>() + SM<900>());

using POTRF = decltype(Base() + TileSize<Base::suggested_tile_size>());
```

--------------------------------

### execute

Source: https://docs.nvidia.com/cuda/cusolverdx/get_started/functions/trsm.html

Device function to perform triangular matrix-matrix solve operations.

```APIDOC
## __device__ void execute(const data_type* A, data_type* B)

### Description
Solves a triangular linear system with multiple right-hand sides using default leading dimensions.

### Parameters
- **A** (const data_type*) - Input batched triangular matrix.
- **B** (data_type*) - Input batched right-hand side matrix, overwritten by the result X on exit.

## __device__ void execute(const data_type* A, const unsigned int lda, data_type* B)

### Description
Solves a triangular linear system specifying the leading dimension of A.

### Parameters
- **A** (const data_type*) - Input batched triangular matrix.
- **lda** (const unsigned int) - Leading dimension of A.
- **B** (data_type*) - Input batched right-hand side matrix, overwritten by the result X on exit.

## __device__ void execute(const data_type* A, data_type* B, const unsigned int ldb)

### Description
Solves a triangular linear system specifying the leading dimension of B.

### Parameters
- **A** (const data_type*) - Input batched triangular matrix.
- **B** (data_type*) - Input batched right-hand side matrix, overwritten by the result X on exit.
- **ldb** (const unsigned int) - Leading dimension of B.

## __device__ void execute(const data_type* A, const unsigned int lda, data_type* B, const unsigned int ldb)

### Description
Solves a triangular linear system specifying leading dimensions for both A and B.

### Parameters
- **A** (const data_type*) - Input batched triangular matrix.
- **lda** (const unsigned int) - Leading dimension of A.
- **B** (data_type*) - Input batched right-hand side matrix, overwritten by the result X on exit.
- **ldb** (const unsigned int) - Leading dimension of B.
```

--------------------------------

### Check complete solver execution

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/description_traits.html

Verifies if a solver descriptor is both complete and configured for execution, enabling the use of Execution Methods.

```cpp
// bool
cusolverdx::is_complete_solver_execution<Solver>::value
cusolverdx::is_complete_solver_execution_v<Solver>
```

--------------------------------

### Retrieve Side Configuration

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/execution_traits.html

Retrieves the side configuration for matrix multiplication.

```cpp
// cusolverdx::side
Solver::side
```

--------------------------------

### POTRF Device Execution

Source: https://docs.nvidia.com/cuda/cusolverdx/get_started/functions/potrf.html

Standard device function signatures for performing batched Cholesky factorization.

```APIDOC
## __device__ void execute(data_type* A, status_type* info)

### Description
Computes batched Cholesky factorization of a Hermitian positive-definite matrix A.

### Parameters
- **A** (data_type*) - Required - Pointer to the batched M×M Hermitian matrix.
- **info** (status_type*) - Required - Pointer to the status array where info[batch_id] indicates the smallest leading minor that is not positive definite if factorization fails.

## __device__ void execute(data_type* A, const unsigned int lda, status_type* info)

### Description
Computes batched Cholesky factorization of a Hermitian positive-definite matrix A with a specified leading dimension.

### Parameters
- **A** (data_type*) - Required - Pointer to the batched M×M Hermitian matrix.
- **lda** (unsigned int) - Required - The leading dimension of the matrix A.
- **info** (status_type*) - Required - Pointer to the status array.
```

--------------------------------

### Initialize Block Operator

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/execution_ops.html

Basic syntax for invoking the Block operator.

```cpp
cusolverdx::Block()
```

--------------------------------

### make_shared_storage_calculator

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_shared.html

Creates an instance of the shared_storage_calculator to compute required shared memory size.

```APIDOC
## shared_storage_calculator make_shared_storage_calculator()

### Description
Creates a new instance of the shared storage calculator used to determine the total shared memory required for various matrices and arrays.

### Signature
`shared_storage_calculator make_shared_storage_calculator();`
```

--------------------------------

### execute

Source: https://docs.nvidia.com/cuda/cusolverdx/get_started/functions/modified_lu.html

Performs a batched modified LU factorization on an orthonormal matrix Q, overwriting it with L and U factors and storing the diagonal sign matrix in S.

```APIDOC
## __device__ void execute(data_type* Q, data_type* S)

### Description
Computes a batched modified LU factorization of an orthonormal matrix Q. The matrix Q is overwritten with the resulting L and U factors, and S is populated with the diagonal sign matrix.

### Parameters
- **Q** (data_type*) - Required - Orthonormal matrix of size M×N per batch. Overwritten with L and U on output.
- **S** (data_type*) - Required - 1D array of size min(M,N) per batch. Contains the diagonal sign matrix on output.

## __device__ void execute(data_type* Q, const unsigned int ldq, data_type* S)

### Description
Computes a batched modified LU factorization of an orthonormal matrix Q with a specified leading dimension.

### Parameters
- **Q** (data_type*) - Required - Orthonormal matrix of size M×N per batch. Overwritten with L and U on output.
- **ldq** (const unsigned int) - Required - Leading dimension of Q.
- **S** (data_type*) - Required - 1D array of size min(M,N) per batch. Contains the diagonal sign matrix on output.
```

--------------------------------

### Basic copy_2d Usage

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_methods.html

Demonstrates transferring matrix data between global and shared memory using operation-based configuration.

```cpp
# include "cusolverdx_io.hpp"

// Example 1: Using operation-based configuration of copy_2d functions
// copy from the global memory to shared memory
cusolverdx::copy_2d<Solver, M, N, cusolverdx::arrangement_of_v_a<Solver>, Solver::BatchesPerBlock>(Ag, lda_g, As, lda_s);
__syncthreads();

// run the solver functions
Solver::execute(As, lda_s, ...);

// copy from the shared memory to the global memory
__syncthreads();
cusolverdx::copy_2d<Solver, M, N, cusolverdx::arrangement_of_v_a<Solver>, Solver::BatchesPerBlock>(As, lda_s, Ag, lda_g);
```

--------------------------------

### Retrieve Suggested Block Dimensions

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/other_traits.html

Accesses the recommended number of threads for a Solver description, which is used if the BlockDim operator is not explicitly defined.

```cpp
// dim3(X, Y==1, Z==1)
Solver::suggested_block_dim
```

--------------------------------

### Set MathDx Root via Command Line

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/installation.html

Configure the project by setting the mathdx_ROOT variable during the CMake configuration step.

```bash
cmake -Dmathdx_ROOT="<your_directory>/nvidia/mathdx/yy.mm/" (...)
```

--------------------------------

### TRSM Device Execution Methods

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/trsm.html

Available device function signatures for executing TRSM, including variants with and without runtime leading dimensions.

```cpp
__device__ void execute(const data_type* A, data_type* B);
// with runtime leading dimensions
__device__ void execute(const data_type* A, const unsigned int lda,
                              data_type* B);
__device__ void execute(const data_type* A,
                              data_type* B, const unsigned int ldb);
__device__ void execute(const data_type* A, const unsigned int lda,
                              data_type* B, const unsigned int ldb);
```

--------------------------------

### cuSolverDx gesvd Device Execution Methods

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/gesvd.html

Overloaded device functions for computing singular values and optional singular vectors, with support for runtime leading dimensions.

```cpp
// Compute singular values only
__device__ void execute(data_type* A, precision_type* S, data_type* workspace, status_type* info);

// Compute singular values only with runtime leading dimension
__device__ void execute(data_type* A, const unsigned int lda, precision_type* S,
                        data_type* workspace, status_type* info);

// Compute singular values and the left/right singular vectors
__device__ void execute(data_type* A, precision_type* S, data_type* U, data_type* VH,
                        data_type* workspace, status_type* info);

// Compute singular values and the left/right singular vectors with runtime leading dimensions
__device__ void execute(data_type* A, const unsigned int lda, precision_type* S,
                        data_type* U, const unsigned int ldu,
                        data_type* VH, const unsigned int ldvh,
                        data_type* workspace, status_type* info);
```

--------------------------------

### gels::execute

Source: https://docs.nvidia.com/cuda/cusolverdx/get_started/functions/gels.html

Device functions to execute the Generalized Linear System Solver.

```APIDOC
## __device__ void execute(data_type* A, data_type* tau, data_type* B)
## __device__ void execute(data_type* A, data_type* tau, data_type* B, const unsigned int ldb)
## __device__ void execute(data_type* A, const unsigned int lda, data_type* tau, data_type* B)
## __device__ void execute(data_type* A, const unsigned int lda, data_type* tau, data_type* B, const unsigned int ldb)

### Description
Solves overdetermined or underdetermined least square problems (min||op(A)X-B||2) using QR or LQ factorization. The function operates in-place, overwriting B with the solution X.

### Parameters
- **A** (data_type*) - Required - Batched M×N matrix. Overwritten by QR or LQ factorization.
- **tau** (data_type*) - Required - Array of size min(M,N) per batch, representing Householder vectors.
- **B** (data_type*) - Required - Batched right-hand side matrix. Overwritten by the solution X.
- **lda** (unsigned int) - Optional - Leading dimension of matrix A.
- **ldb** (unsigned int) - Optional - Leading dimension of matrix B.
```

--------------------------------

### GETRS Partial Pivot Device Functions

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/getrs.html

Device function signatures for solving linear systems with partial pivoting, requiring an ipiv array for row interchanges.

```cpp
__device__ void execute(const data_type* A, const int* ipiv, data_type* B);
// with runtime leading dimensions
__device__ void execute(const data_type* A, const unsigned int lda,
                        const int* ipiv,
                        data_type* B);
__device__ void execute(const data_type* A,
                        const int* ipiv,
                        data_type* B, const unsigned int ldb);
__device__ void execute(const data_type* A, const unsigned int lda,
                        const int* ipiv,
                        data_type* B, const unsigned int ldb);
```

--------------------------------

### Check Solver Support

Source: https://docs.nvidia.com/cuda/cusolverdx/api/other_traits.html

Use is_supported or is_supported_v to verify if a solver operation is supported on a specific CUDA architecture.

```cpp
// true if Solver is supported on the provided CUDA architecture
cusolverdx::is_supported<Solver, Arch>::value;
cusolverdx::is_supported_v<Solver, Arch>;
```

--------------------------------

### Solver::execute method signature

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/methods.html

Computes singular values and optional singular vectors with runtime leading dimensions.

```cpp
// Compute singular values, and optionally the left and/or right singular vectors
 // with runtime leading dimension of `U` and `VT`
 void Solver::execute(typename Solver::a_data_type* A, const unsigned int runtime_lda,
                    typename Solver::a_precision* sigma,
                    typename Solver::a_data_type* U, const unsigned int ldu,
                    typename Solver::a_data_type* VT, const unsigned int ldvt,
                    typename Solver::a_workspace_type* workspace, typename Solver::status_type* info);
```

--------------------------------

### Initialize Cluster Operator

Source: https://docs.nvidia.com/cuda/cusolverdx/0.5.0/api/execution_ops.html

Basic syntax for invoking the Cluster operator.

```cpp
cusolverdx::Cluster()
```

--------------------------------

### void execute(const data_type* A, const unsigned int lda, data_type* B, const unsigned int ldb)

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/potrs.html

Solves the linear system AX=B with specified leading dimensions for both A and B.

```APIDOC
## __device__ void execute(const data_type* A, const unsigned int lda, data_type* B, const unsigned int ldb)

### Parameters
- **A** (const data_type*) - The batched M×M Hermitian matrix.
- **lda** (const unsigned int) - The leading dimension of matrix A.
- **B** (data_type*) - The batched right-hand side matrix.
- **ldb** (const unsigned int) - The leading dimension of matrix B.
```

--------------------------------

### GESVD Execution Methods

Source: https://docs.nvidia.com/cuda/cusolverdx/api/methods.html

Overloaded methods for performing Singular Value Decomposition on general matrices.

```cpp
void Solver::execute(typename Solver::a_data_type* A, typename Solver::a_precision* sigma,
                      typename Solver::a_data_type* workspace, typename Solver::status_type* info);
```

```cpp
void Solver::execute(typename Solver::a_data_type* A, const unsigned int runtime_lda,
                      typename Solver::a_precision* sigma,
                      typename Solver::a_data_type* workspace, typename Solver::status_type* info);
```

```cpp
void Solver::execute(typename Solver::a_data_type* A, typename Solver::a_precision* sigma,
                      typename Solver::a_data_type* U, typename Solver::a_data_type* VT,
                      typename Solver::a_data_type* workspace, typename Solver::status_type* info);
```

```cpp
void Solver::execute(typename Solver::a_data_type* A, const unsigned int runtime_lda,
                    typename Solver::a_precision* sigma,
                    typename Solver::a_data_type* U, const unsigned int ldu,
                    typename Solver::a_data_type* VT, const unsigned int ldvt,
                    typename Solver::a_data_type* workspace, typename Solver::status_type* info);
```

--------------------------------

### Check complete solver descriptor

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/api/description_traits.html

Verifies if a solver descriptor is complete, meaning it has all necessary description operators applied.

```cpp
// bool
cusolverdx::is_complete_solver<Solver>::value
cusolverdx::is_complete_solver_v<Solver>
```

--------------------------------

### void execute(const data_type* A, data_type* B, const unsigned int ldb)

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/potrs.html

Solves the linear system AX=B with a specified leading dimension for B.

```APIDOC
## __device__ void execute(const data_type* A, data_type* B, const unsigned int ldb)

### Parameters
- **A** (const data_type*) - The batched M×M Hermitian matrix.
- **B** (data_type*) - The batched right-hand side matrix.
- **ldb** (const unsigned int) - The leading dimension of matrix B.
```

--------------------------------

### GETRS No Pivot Device Functions

Source: https://docs.nvidia.com/cuda/cusolverdx/0.4.0/get_started/functions/getrs.html

Device function signatures for solving linear systems without pivoting, including variants for runtime leading dimensions.

```cpp
__device__ void execute(const data_type* A, data_type* B);
// with runtime leading dimensions
__device__ void execute(const data_type* A, const unsigned int lda,
                        data_type* B);
__device__ void execute(const data_type* A,
                        data_type* B, const unsigned int ldb);
__device__ void execute(const data_type* A, const unsigned int lda,
                        data_type* B, const unsigned int ldb);
```
Tokens:
Raw