Calculating the Singular Values of Many Small Matrices on GPUs

Abstract

This paper presents a fast and robust approach to evaluate the singular values of small (e.g., $4\times 4$, $5\times 5$) matrices on single- and multi-Graphics Processing Unit (GPU) systems, enabling the modulation of the accuracy–speed trade-off. Targeting applications that require only computations of the SVs in electromagnetics (e.g., Multiple Input Multiple Output—MIMO link capacity optimization) and emerging deep-learning kernels, our method contrasts with existing GPU singular value decomposition (SVD) routines by computing singular values only, thereby reducing overhead compared to full-SVD libraries such as cuSOLVER’s gesvd and MKL’s desvg. The method uses four steps: interlaced storage of the matrices in GPU global memory, bidiagonalization via Householder transformations, symmetric tridiagonalization, and root finding by bisection using Sturm sequences. We implemented the algorithm in CUDA and evaluated it on different single- and multi-GPU systems. The approach is particularly suited for the analysis and design of multiple-input/multiple-output (MIMO) communication links, where thousands of tiny SVDs must be computed rapidly. As an example of the satisfactory performance of our approach, the speed-up reached for large matrix batches against cuSOLVER’s gesvd has been around 20 for $4\times 4$ matrices. Furthermore, near-linear scaling across multi-GPUs systems has been reached, while maintaining root mean square errors below $2.3\times {10}^{-7}$ in single precision and below $2.3\times {10}^{-13}$ in double precision. Tightening the tolerance from $\delta ={10}^{-7}$ to $\delta ={10}^{-9}$ increased the total runtime by only about 10%.

1. Introduction

One of the most useful result of linear algebra is the singular value decomposition (SVD), which is widely employed in many areas of applied research like inverse problems [1], data compression [2,3], and machine learning [4,5]. In electromagnetics, it has been fruitfully used in, among other topics, near-field antenna sampling [6], imaging [7,8], and optimization of multiple-input/multiple-output (MIMO) systems [9].

1.1. Application Context

In some of the applications, the full SVD computation, namely singular vectors and singular values (SVs), is required, while in others only SVs are necessary. For example, MIMO is a wireless communication technology that uses multiple antennas working in both transmission and reception modalities to increase the channel capacity over that of a conventional single-input/single-output (SISO) system. In this case, the channel capacity is a crucial parameter to evaluate the performance of a communication system, and the problem arises of how to define the most convenient antenna configuration to maximize it, given the signal-to-noise ratio and the propagation environment. The channel capacity is related to the SVs of the channel matrix, its optimization requires repeated calculation of the SVs [9], and so accelerating their computation is crucial even in the case of classically sized (i.e., $4\times 4$, $5\times 5$) MIMO [10]. Of course, when only the SVs are required, calculating the SVs alone instead of the full SVD turns out to be more convenient, this problem being computationally less burdened.

Other applications requiring the computation of large batches of tiny matrices are astrophysics [11] and deep learning [12,13]. For example, in an astrophysical ordinary differential equation (ODE) solver, multiple subregions of the domain of interest are simulated, and each subregion corresponds to a small linear system to be analyzed. In the framework of deep learning, the singular values of the network layers are computed [12] in order to be, for example, bounded [13] to prevent the effect of ill-conditioning during training. Speaking more generally, a recent trend is occurring in redefining linear algebra algorithms for batches of small matrices [14]. For example, it has been shown that randomly sampling a matrix can lead to satisfactory approximations of its first singular values [15].

Performing an accurate and fast computation of SVs is a sensitive issue, especially when the calculation is part of a very large number of iterations, as in optimization problems [6,7,8,9]. In fact, in addition to the repeated evaluation, each iteration can require the computation of the SVs of many different matrices, either for global or local optimizations. Graphics Processing Units (GPUs) are a candidate computing platform to accelerate this task [16,17,18].

1.2. Limitations of Literature

The computation of singular values only on conventional computational platforms has been extensively explored in the literature [19]. Modern approaches to this problem have been presented in [20], where the QR method for upper-triangular semiseparable matrices has been applied after the original matrix has been reduced to bidiagonal form, and in [21], where a deep neural network-based approach has been proposed. Unfortunately, the problem dealt with in this paper has received little attention, since the approaches available for GPU processing have mainly focused on the computation of the full SVD of large batches of large matrices. The few GPU algorithms dealing with large batches of small/medium-sized matrices have tackled the problem of the full or reduced SVD calculation involving the computation of both singular values and singular vectors, and not the singular values only. For example, in [22], a W-cycle Jacobi SVD on an NVIDIA V100 running $4.5$ faster than cuSOLVER has been proposed for the SVD computation of small/moderate matrix sizes (e.g., $32\times 32$, $48\times 48$). In [23], a throughput exceeding $800GFLOP/s$ in single-precision arithmetic using register/shared-memory optimizations has been achieved on K40/P100 systems for the SVD calculation of medium-sized matrices (e.g., $512\times 512$). Finally, in [14], batched linear algebra solutions, including SVD, were investigated for matrices from small/moderate to medium sizes (up to $512\times 512$).

1.3. Our Contributions

In this paper, we develop a fast and robust approach to evaluate the SVs of small (e.g., $4\times 4$, $5\times 5$) matrices on GPUs capable of modulating the accuracy-speed trade-off. The primary requirement is developing GPU routines having a convenient computational cost so as to be faster than state-of-the-art libraries running on GPUs.

Obviously, the achievable effectiveness and efficiency of the approach are related to the features of the involved matrices that can be exploited as a priori knowledge and to the capability of taking advantage of the underlying architecture of the GPU at its best. Concerning the former point, many of the algebraic properties that lead to benefits in the computation of the SVs of large matrices, e.g., symmetry [24] or rotational symmetry [25], have little effect in the case of small matrices. Accordingly, and also to keep the approach general, in this paper we will exploit no such properties. Regarding the latter point, for the problem here considered and to gain the required effectiveness and efficiency, the computational scheme must be properly devised to exploit at best the Single Instruction Multiple Data (SIMD) paradigm of GPUs. We note that, although many schemes for the SV calculation have been developed, choosing and implementing the best approach matching the SIMD pattern is by no means an easy task when a large number of small matrices is involved. Here, for the very first time, we present an approach for the calculation of the SVs of many small matrices based on four steps that have been properly selected to match the SIMD architecture of a GPU:

• Interlaced storage of the matrices in GPU global memory;
• Bidiagonalization using Householder matrices [26];
• Symmetric tridiagonalization of the bidiagonal matrix;
• Root finding by bisection using Sturm sequences [27].

In summary, the contributions of this paper are threefold:

• we introduce a novel, multi-GPU algorithm for computing the singular values of many small matrices via interlaced global memory storage, Householder-based bidiagonalization, tridiagonalization, and Sturm-sequence bisection, allowing direct control of the accuracy–speed trade-off via the tolerance parameter $\delta$;
• we provide a CUDA implementation that achieves remarkable speedups over other existing approaches and whose target accuracy can be set to match MATLAB’s state-of-the-art svd routine;
• we demonstrate the algorithm scalability across multiple GPUs and show the computational performance behavior against the target accuracy.

The implementation has been devised as follows.

First, a proper interlaced storage of the data matrices has been set up to guarantee coalesced memory accesses. This has been a key point of implementation, as the algorithm is memory bound. Second, it should be noted that the bidiagonalization step is not easy to parallelize and that the matrices involved are small, which makes thread cooperation not convenient for this task. Accordingly, in the proposed computational scheme, each GPU thread is in charge of calculating the whole bidiagonalization for a different matrix. We stress that this choice is strictly related to the size of the relevant matrices and that, conversely, for large matrices such a strategy would lead to unsatisfactory results. Third, since the approach based on Sturm sequences turns root finding into an embarrassingly parallel task, each different matrix eigenvalue is carried out in parallel by a different GPU thread, so that different GPU threads can compute eigenvalues of the same matrix in parallel.

1.4. Differences with the Literature

Targeting applications requesting the calculation of the SVs only, compared to the approaches in [14,22,23], our computational scheme calculates singular values alone, without the singular vectors, which eases the burden and enables much higher speedups against standard library routines compared to [14,22,23]. Furthermore, as will be shown, by a tunable tolerance, our approach is capable of reaching the same singular value computational accuracy of MKL or MATLAB, while the precision issue has been neglected in [22], can remain undefined in [23] in the cases of lack of convergence, is implicitly inherited by LAPACK’s Householder kernels, and is not directly controlled in [14].

It should also be noticed that, in recent years, the availability of multi-GPU systems, single nodes or multi-node, has spread [28]. In order to take into account such an availability, another key point of the developed approach has been exploiting part or all of the GPUs of the computing system by subdividing the batch of matrices for which the SVs are computed across the different underlying devices.

We finally want to stress that, to the best of our knowledge, the approach proposed for the multi-GPU calculation of the SVs of a large number of small matrices is an absolute novelty. Among others [14,23], one of the operational GPU routines available for the calculation of SVs of matrix batches is gesvd of the cuSOLVER library [29]. Nevertheless, as we shall see, the performance of such a routine is worse than the proposed computational scheme. Furthermore, throughout the literature, attention has been almost exclusively paid to GPU computations of the SVD for large matrices [30,31].

The algorithm has been implemented in the CUDA language and tested on several single- and multi-GPU systems.

The original, preliminary approach in [10] has been thoroughly optimized. Also, at variance with [10], the performance has been investigated over different single- and multi-GPU architectures, also in comparison to other GPU-based SV calculation routines.

1.5. Paper Layout

The paper is organized as in Figure 1. In Section 2, the most common methods to compute the SV are discussed, and the proposed approach is presented. In Section 3, the strategy adopted to implement the algorithm on the GPU is described. Section 4 shows the numerical results. Finally, in Section 5, conclusions are drawn and future developments are foreseen. A detailed description of the bidiagonalization (Algorithm 1) is given as well as CUDA-like snippets for the memory reorganization (Algorithm 2), bidiagonalization (Algorithm 3), bisection (Algorithm 4) and routines used in the root finding step (Algorithms A1 and A2) are reported. A list of all abbreviations used throughout the manuscript is finally provided in the Appendix A.

2. The Approach

2.1. Statement of the Problem

In the following, we will deal with the particular case of real matrices, the extension to complex matrices being a matter of a future activity. Furthermore, we will use the same MATLAB-like notation of [26] for vectors, matrices, and operations among them, which can now be considered as a standard in linear algebra applications. Finally, with no loss of generality, we will consider matrices $A\in {\mathbb{R}}^{mxn}$, with $m\le n$. Obviously, in the case $m>n$, it will be sufficient to compute the singular values of $A^T$, the superscript T denoting transposition, since they coincide with those of A.

The problem of computing the singular values of a matrix can be theoretically reduced to that of finding the eigenvalues of a symmetric matrix. Indeed, the singular values of A are the square roots of the eigenvalues of the matrix $A^{T}A$, or, equivalently, $A{A}^T$. Likewise, the left (or right) singular vectors of A are equal to the eigenvectors of $A{A}^T$ (or $A^{T}A$).

2.2. Literature Review

Much literature is available for the full computation (i.e., both singular values and singular vectors) of the SVD of a matrix. The two most popular techniques are

• the Golub-Reinsch algorithm [26,32];
• the one-sided Jacobi algorithm [33,34].

Other techniques include the bisection [27] and the divide and conquer [35] approaches.

2.2.1. The Golub-Reinsch Algorithm

The Golub–Reinsch algorithm is the standard for the computation of the SVD, and it is also used in the LAPACK package [26,32]. In particular, the Golub–Reinsch algorithm is a two-step one. The input matrix is first reduced to a bidiagonal matrix using a series of Householder transformations and then the bidiagonal matrix is diagonalized by performing implicit QR iterations.

2.2.2. The One-Sided Jacobi Method

The one-sided Jacobi method, instead, makes use of a sequence of pair-wise column orthogonalizations such that the input matrix is replaced, after a proper number of iterations, with a sufficiently orthogonal matrix according to a given tolerance [26,32].

2.2.3. Differences Between the Golub-Reinsch Algorithm and the One-Sided Jacobi Method

The Golub-Reinsch method appears to be more difficult to be parallelized, and its sequential version is faster than Jacobi-based algorithms on sequential machines [36,37]. On the other hand, the Jacobi method exhibits some degree of parallelism [38,39,40] and proves a better relative accuracy [41]. However, neither of the two methods cited above appears to be amenable to a meaningful parallel implementation with tiny (e.g., $4\times 4$, $5\times 5$) matrices. Furthermore, we are here interested in the calculation of the SVs only, while the two above approaches aim at the full SVD computation.

2.2.4. Outline of Our Approach

As mentioned in the Introduction, the proposed approach to face the problem of computing the SVs of many small matrices is a three-step one. In Figure 2, we outline the four major stages of our algorithm:

• Interlaced memory storage. A pre-processing step amounts to interlace the different matrices to be processed in order to enable coalesced memory accesses.
• Bidiagonalization. The first step consists of reducing the input matrix into a bidiagonal form, following the same lines as the Golub-Reinsch algorithm.
• Tridiagonalization. The second step performs a symmetric tridiagonalization of the bidiagonal matrix.
• Root finding. The third step searches for the roots of the characteristic polynomial associated with the symmetric tridiagonal matrix by bisection using Sturm sequences.

In the following sections, we will briefly recall only the main mathematical principles behind the approach necessary to describe the implementation.

Figure 2. Block diagram of the proposed approach.

Figure 2. Block diagram of the proposed approach.

2.3. First Step: Bidiagonalization

In this step, the generic real matrix $A\in {\mathbb{R}}^{mxn}$ is bidiagonalized, namely, it is decomposed as

$$A=PB{Q}^T,$$

(1)

where $B\in {\mathbb{R}}^{mxn}$ is an upper bidiagonal matrix and $P\in {\mathbb{R}}^{mxm}$ and $Q\in {\mathbb{R}}^{nxn}$ are orthogonal Householder matrices. The bidiagonalization is achieved by applying a series of Householder transformations to the matrix A with the aim of zeroing the elements below the diagonal along with the elements above the first superdiagonal. The bidiagonalization algorithm that we have employed and that we are now going to briefly recall is that presented by Golub and Van Loan in [26]. It is an iterative algorithm having a finite number of iterations. In particular, the number of iterations is $n-1$ if $m=n$; otherwise, the number of iterations is n.

2.3.1. Householder Reflection

Before describing the bidiagonalization algorithm, we recall that a possible Householder vector h to eliminate entries below the top entry, namely, to make a given vector $x\in {\mathbb{R}}^{kx1}$ proportional to a vector $e\in {\mathbb{R}}^{kx1}={(1,0,\dots ,0)}^T$ is given by [26]

$$h=x+\alpha e,$$

(2)

where $\alpha =\pm \parallel x\parallel$. According to Equation (2) and on defining $\beta =2/\left(h^{T}h\right)$ and $I_k$ as the identity matrix of dimensions $k\times k$, then [26]

$$(I_k-\beta h{h}^T)x=\alpha e.$$

(3)

The vector h is also known as a reflector. A stable way to compute the Householder reflectors is described in [26] (see Algorithm 5.1.1, p. 210).

We can think of bidiagonalization as alternately “zeroing-out” columns and rows with simple reflection operations. Each reflector acts like placing a mirror at an angle to “bounce” a column (or row) so that all but its first entry vanish (Figure 3).

2.3.2. Iterations

The first operation of iteration #1 consists of selecting a left Householder vector $u^{\left(1\right)}$ of length m to zero the elements of A below the first element of the first column. The vector x is represented in this case by $A(1:m,1)$. Using $u^{\left(1\right)}$ as a reflector, we define the matrix $A_1^{\prime }$ as [26]

$$A_1^{\prime }=(I_m-{\beta }_u^{\left(1\right)}{u}^{\left(1\right)}{u^{\left(1\right)}}^T)A=\left[\begin{array}{ccccc}X& X& X& \dots & X\\ 0& X& X& \dots & X\\ 0& X& X& \dots & X\\ ⋮& X& X& \dots & X\\ 0& X& X& \dots & X\end{array}\right].$$

(4)

The second operation consists of defining a new matrix of size $m\times (n-1)$ as $A_1^{″}=A_1^{\prime }(1:m,2:n)$, selecting a right Householder vector $v^{\left(1\right)}$ of length $n-1$ to zero the elements of $A_1^{″}$ to the right of the first element of the first row. In this case, the vector x is represented by $A_1^{″}(1,1:n-1)$. It should be noticed that a further requirement is that the Householder vectors be normalized such as $v^{\left(1\right)}\left(1\right)=1$. Using $v^{\left(1\right)}$ as a reflector, we define the matrix $A_2$ as [26]

$$A_2=A_1^{″}(I_{n-1}-{\beta }_v^{\left(1\right)}{v}^{\left(1\right)}{v^{\left(1\right)}}^T)=\left[\begin{array}{ccccc}x& x& 0& \dots & 0\\ 0& x& x& \dots & x\\ 0& x& x& \dots & x\\ ⋮& x& x& \dots & x\\ 0& x& x& \dots & x\end{array}\right]$$

(5)

where ${\beta }_u^{\left(1\right)}=2/({u^{\left(1\right)}}^T{u}^{\left(1\right)})$, ${\beta }_v^{\left(1\right)}=2/({v^{\left(1\right)}}^T{v}^{\left(1\right)})$ and I is the identity matrix. In this way, $A_2$ shows zeros below the diagonal in the first column and to the right of the superdiagonal element of the first row, as requested.

The algorithm then proceeds in an analogous way with the annihilation of the second column $A(2:m,2)$ and row $A(2,3:n)$ and so on. If $m>n$, there are n columns and $n-2$ rows to eliminate. At the end of this process, an upper bidiagonal matrix B is obtained such that [26]

$$P=\prod _{i=1}^n{P}_i\mathrm{and}Q=\prod _{i=1}^{n-2}{Q}_i$$

(6)

with

$$P_i=(I-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{u^{\left(i\right)}}^T)\mathrm{and}{Q}_i=(I-{\beta }_v^{\left(i\right)}{v}^{\left(i\right)}{v^{\left(i\right)}}^T).$$

(7)

The left Householder vectors $u^{\left(i\right)}$’s are vectors of length m with $i-1$ leading zeros, while the right Householder vectors $v^{\left(i\right)}$’s are vectors of length n with i leading zeros. They are needed to zero the elements of the ith column below the diagonal as well as to zero the elements of the ith row to the right of the superdiagonal element, respectively. Moreover, both left and right Householder vectors are normalized such that $u^{\left(i\right)}\left(i\right)=1$ and $v^{\left(i\right)}(i+1)=1$.

From the computational standpoint, each left or right update of the bidiagonalization process consists of a Matrix-Vector Multiplication (MVM) and a rank-one update. Indeed, looking at the left matrix update at the ith step, the involved algebraic operations can be arranged as

$$(I-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{u^{\left(i\right)}}^T)A_i=\underset{\mathrm{R}\mathrm{a}\mathrm{n}\mathrm{k}\text{-}\mathrm{o}\mathrm{n}\mathrm{e} \mathrm{u}\mathrm{p}\mathrm{d}\mathrm{a}\mathrm{t}\mathrm{e}}{\underbrace{A_i-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}\underset{\mathrm{M}\mathrm{V}\mathrm{M}}{\underbrace{{y^{\left(i\right)}}^T}}}}$$

(8)

where ${y^{\left(i\right)}}^T\in {\mathbb{R}}^{1xn}={u^{\left(i\right)}}^T{A}_i$ is the relevant MVM and $-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{y^{\left(i\right)}}^T$ is the rank-one matrix involved in the update. The same holds true for the right matrix updates. Accordingly, each bidiagonalization step costs two MVMs and two rank-one updates. Actually, the ith bidiagonalization step can be written as

$$A_{i+1}=(I-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{u^{\left(i\right)}}^T)A_i(I-{\beta }_v^{\left(i\right)}{v}^{\left(i\right)}{v^{\left(i\right)}}^T).$$

(9)

It should be noticed that the operations in Equation (9) must be executed in a prescribed order since the ith right Householder vector $v^{\left(i\right)}$ is known only after the ith column has been annihilated. In other words, $v^{\left(i\right)}$ can be computed only after the ith row of $(I-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{u^{\left(i\right)}}^T)A_i$ has been calculated. Nevertheless, let us ignore this issue for a while. With a little algebra, Equation (9) can be recast as

$$A_{i+1}=\underset{\mathrm{R}\mathrm{a}\mathrm{n}\mathrm{k}\text{-}\mathrm{t}\mathrm{w}\mathrm{o} \mathrm{u}\mathrm{p}\mathrm{d}\mathrm{a}\mathrm{t}\mathrm{e}}{\underbrace{A_i-u^{\left(i\right)}{z}^{\left(i\right)}-\underset{\mathrm{MVM}}{\underbrace{w^{\left(i\right)}}}{v^{\left(i\right)}}^T}}$$

(10)

where

$$w^{\left(i\right)} = {\beta }_v^{\left(i\right)}{A}_i{v}^{\left(i\right)},$$

(11)

$$x^{\left(i\right)} = {\beta }_u^{\left(i\right)}{A}_i^T{u}^{\left(i\right)},$$

(12)

$$\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}{z}^{\left(i\right)} = \underset{\mathrm{M}\mathrm{V}\mathrm{M}}{\underbrace{{x^{\left(i\right)}}^T}}-{\beta }_v^{\left(i\right)}\left({x^{\left(i\right)}}^T{v}^{\left(i\right)}\right){v^{\left(i\right)}}^T.$$

(13)

Assuming we perform an in-place update of the input matrix, Equations (10)–(13) allow the set up of the bidiagonalization step as two MVMs and a single rank-two update, where $u^{\left(i\right)}{z}^{\left(i\right)}-w^{\left(i\right)}{v^{\left(i\right)}}^T$ is the rank-two matrix involved in such an update. In other words, the mentioned equations enable deferring one of the two left rank-one matrix updates and saving a read-write operation as compared to the scheme in Equation (8). Concerning now the problem of knowing $v^{\left(i\right)}$ in advance, this issue can be solved by updating at least the ith row of the matrix $A_i$ prior to applying Equations (10)–(13). The described bidiagonalization approach is summarized in Algorithm 1.

Algorithm 1 Bidiagonalization.

if  $m==n$  then    $numiter=n-1$; else    $numiter=n$; end for  $i=1:n-2$  do  1. Compute the ith left Householder vector $u^{\left(i\right)}$ relative to the column $A_i(i:m,i)$.  2. Update the ith row of $A_i$ to compute $v^{\left(i\right)}=A_i(i,i+1:n)-{\beta }_u^{\left(i\right)}{u^{\left(i\right)}}^T{A}_i(i:m,i+1:n)$.  3. Compute the ith right Householder vector $v^{\left(i\right)}$ from ${\widehat{v}}^{\left(i\right)}$.  4. Compute $w^{\left(i\right)}\mathrm{and}{x}^{\left(i\right)}$ from (11) and (12), respectively.  5. Compute $z^{\left(i\right)}$ from (13).  6. Update the matrix $A_i$ using (10). end for  $i=n-2:numiter$  do  1. Compute the ith left Householder vector $u^{\left(i\right)}$ relative to the column $A_i(i:m,i)$.  2. Update the input matrix using (8):   $A_{i+1}=(I-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{u^{\left(i\right)}}^T)A_i=A_i-{\beta }_u^{\left(i\right)}{u}^{\left(i\right)}{y^{\left(i\right)}}^T$. end

2.4. Second Step: Tridiagonalization

Once the input matrix A has been transformed into a bidiagonal form B, then the singular values of A could be obtained by applying an implicit-shift QR algorithm to the symmetric tridiagonal matrix $T=B^{T}B$.

It is known that, for numerical reasons, the explicit formation of the matrix T should be avoided since, depending on the nature and on the algebraic properties of B, the explicit computation of T could result in unnecessary loss of accuracy, especially in the calculation of the smallest singular values [27,32]. To avoid this problem, the Golub-Reinsch method applies a series of implicit QR factorizations directly to the matrix B; see [19,26,32,42].

However, one of the major drawbacks of the QR algorithm is that it does not show enough parallelism since each orthogonal similarity transformation needed to reduce the bidiagonal matrix B to a diagonal form is dependent upon the previous one. In principle, this problem could be overcome by letting each thread compute the QR algorithm for a different matrix. Indeed, we recall that our aim is to calculate the singular values of a large number of small matrices so that the computational load of each thread would be light. Unfortunately, such an approach would not be convenient on GPUs for two reasons. First, each matrix could require a quite different number of iterations to converge, thus resulting in a highly unbalanced load among the threads and so failing to keep all the GPU multiprocessors busy. Moreover, individual threads of commonly available multi-core CPUs would execute each sequential orthogonal similarity transformation in a shorter time than a GPU thread.

An alternative approach to the Golub-Reinsch algorithm is represented by the one-sided Jacobi method, which was (for much of the time) neglected due to its slowness as compared to the standard bidiagonalization-based methods. Opposite to that, it has aroused much attention in the last few years thanks to recent developments in [36,37,41].

In this paper, after reducing the input matrix A to a bidiagonal matrix B, B is transformed to a tridiagonal matrix $T=B^{T}B$ and the symmetric tridiagonal eigenvalue problem for the matrix T is solved by the method of root finding by bisection using Sturm sequences described in the next Subsection. In this way, the singular values of A can be simply determined as the square roots of the eigenvalues of T.

As already underlined, forming matrix T is not the best choice from the standpoint of the numerical accuracy and may introduce non-negligible errors in the smallest singular values. Nevertheless, the developed method is meant for those applications for which the smallest singular values have a negligible role. Examples of those applications are the already mentioned singular value optimization problems arising in antenna near-field sampling [6], imaging [7,8] and MIMO system optimization [9].

2.5. Third Step: Root Finding

In order to describe the method of root finding by bisection using Sturm sequences to find the eigenvalues of T, we assume that the matrix T has the following form [27]

$$T=\left[\begin{array}{ccccc}{\alpha }_1& {\beta }_1& 0& \dots & 0\\ {\beta }_1& {\alpha }_2& {\beta }_2& \dots & 0\\ 0& \ddots & \ddots & \dots & 0\\ ⋮& \ddots & \ddots & \ddots & {\beta }_{n-1}\\ 0& \dots & \dots & {\beta }_{n-1}& {\alpha }_n\end{array}\right].$$

(14)

We also assume that ${\beta }_i\ne 0\forall i$, otherwise the problem dimensions can be reduced.

Let us consider at this point the so-called Sturm polynomials sequence

$$\{p_0\left(x\right),p_1\left(x\right),\dots ,p_n\left(x\right)\},$$

(15)

whose polynomials are defined as [42,43].

$$p_0\left(x\right) = 1$$

(16a)

$$\hspace{1em}\hspace{1em}{p}_1\left(x\right) = {\alpha }_1-x$$

(16b)

$$\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}{p}_i\left(x\right) = ({\alpha }_i-x)p_{i-1}\left(x\right)-{\beta }_{i-1}^2{p}_{i-2}\left(x\right),i=2,\dots ,n$$

(16c)

In (16), $p_n\left(x\right)$ corresponds to the characteristic polynomial associated with T so that its roots coincide with the eigenvalues of T.

The Sturm sequence has the fundamental property that the number of disagreements in sign between consecutive evaluations of the polynomials for $i=0,\dots ,n$ at a certain point x, say $S\left(x\right)$, coincides with the number of eigenvalues of T less than x [42]. Such a property is very helpful in bracketing the interval onto which searching for the eigenvalues of T by bisection. Of course, to start a bisection procedure and achieve fast convergence, a good starting interval is required. This can be achieved by aid of the Gerschgoring circle theorem [26], which states that, for a complex-valued matrix $M\in {\mathbb{C}}^{nxn}$, the spectrum $\sigma \left(M\right)$ is included in the union of the Gershgorin circles $R_i$, with $i=1,\dots ,n$, where $R_i$ is defined as

$$R_i=\{z\in \mathbb{C}:|z-m_{ii}|\le \sum _{\stackrel{j=1}{j\ne i}}^n|m_{ij}\left|\right\}$$

(17)

and where the elements of M have been denoted by $m_{ij}$. Of course, when the matrix T is real, the Gerschgoring circles $R_i$ become segments. In addition, taking into account that in the specific case $M=T$ is not only real but also symmetric tridiagonal, then computing the endpoints of the interval $I=[x_{min},x_{max}]$ containing all the eigenvalues of T follows immediately from (17), namely [44].

$$\begin{array}{cc}\hfill x_{min}& =\underset{1\le i\le n}{min}[{\alpha }_i-\left(\right|{\beta }_{i-1}|+|{\beta }_i\left|\right)]\hfill \end{array}$$

(18a)

$$\begin{array}{cc}\hfill x_{max}& =\underset{1\le i\le n}{max}[{\alpha }_i+\left(\right|{\beta }_{i-1}|+|{\beta }_i\left|\right)]\hfill \end{array}$$

(18b)

We now want to show how, using the bisection algorithm and the above mentioned property of the Sturm sequences, it is possible to find the eigenvalues according to their indices, remembering also that they are ordered in a descending order. Let us calculate the ith eigenvalue of T in $(x_{min},x_{max})$. Let also a and b be the endpoints of the current search interval and c the middle point of the interval $(a,b)$. The algorithm starts by assigning

$$a^{\left(0\right)} = x_{min}$$

(19a)

$$b^{\left(0\right)} = x_{max}$$

(19b)

$$\hspace{1em}\hspace{1em}\hspace{1em}\hspace{1em}{c}^{\left(0\right)} = (x_{min}+x_{max})/2$$

(19c)

and checks whether the ith eigenvalue $s_i$ belongs to $(a^{\left(0\right)},c^{\left(0\right)})$ or to $(c^{\left(0\right)},b^{\left(0\right)})$. This can be carried out by evaluating $S\left(c^{\left(0\right)}\right)$, namely, the disagreements in sign of consecutive polynomials of the Sturm sequence at $c^{\left(0\right)}$, so that from the number of changes of signs of such polynomials the number of eigenvalues of T less than $c^{\left(0\right)}$ can be deduced. Since the eigenvalues of T are n, then $n-S\left(c^{\left(0\right)}\right)$ coincides with the number of eigenvalues larger than $c^{\left(0\right)}$. Then, if $n-S\left(c^{\left(0\right)}\right)<i$, this means that $s_i\in (a^{\left(0\right)},c^{\left(0\right)})$. Accordingly, the algorithm sets

$$\begin{array}{cc}\hfill a^{\left(1\right)}& =x_{min}\hfill \end{array}$$

(20a)

$$\begin{array}{cc}\hfill b^{\left(1\right)}& =c^{\left(0\right)}\hfill \end{array}$$

(20b)

or

$$\begin{array}{cc}\hfill a^{\left(1\right)}& =c^{\left(0\right)}\hfill \end{array}$$

(21a)

$$\begin{array}{cc}\hfill b^{\left(1\right)}& =x_{max}\hfill \end{array}$$

(21b)

otherwise. This process can be straightforwardly generalized. At the kth step, the algorithm operates as follows

$$\left\{\begin{array}{cc}\mathrm{if}S\left(c^{\left(k\right)}\right)>n-i,\hfill & \mathrm{then}{b}^{(k+1)}=c^{\left(k\right)}\mathrm{and}{a}^{(k+1)}=a^{\left(k\right)}\hfill \\ \mathrm{if}S\left(c^{\left(k\right)}\right)\le n-i,\hfill & \mathrm{then}{a}^{(k+1)}=c^{\left(k\right)}\mathrm{and}{b}^{(k+1)}=b^{\left(k\right)}\hfill \end{array}\right.$$

(22)

The iterations proceed until the search interval becomes small enough, up to a given tolerance. In particular, the stop condition is

$$|b^{\left(k\right)}-a^{\left(k\right)}|\le \delta (|a^{\left(k\right)}|+|b^{\left(k\right)}\left|\right)$$

(23)

where $\delta$ is chosen so as to tune the accuracy of the algorithm against the computational burden.

This method turns out to be embarrassingly parallel since each eigenvalue of each matrix can be evaluated independently by each GPU thread.

Particular attention should be, however, paid to the calculation of $S\left(x\right)$, which is, in this paper, performed through the classical function Count(x) [27,43]. A discussion on a stable and accurate implementation of the Count(x) function is reported in the Appendix A.

3. GPU Implementation

In this section, we provide details of the implementation of our approach in a numerical code written in CUDA language for the parallel execution on a GPU. In particular, we account for five different key points of the implementation:

• data organization in device global memory to foster coalesced memory accesses;
• parallelization of the bidiagonalization step;
• parallelization of the tridiagonalization step;
• parallelization of the root finding approach using Sturm sequences;
• multi-GPU processing.

The developed code has been heavily optimized and micro-optimized. Among others, and besides the mentioned proper arrangement of the data that will be discussed in the next Subsection, the most useful optimization actions have been the following:

• Use of const and __restrict__ whenever applicable to the input parameters of __global__ functions, which, as known, enables the use of the read-only data cache available since the Kepler architecture [45].
• Use of template kernel parameters, which requires to know relevant problem parameters, as the matrix pool size K and the matrix dimensions m and n, at compile time. We have actually implemented two versions of the approach, one using template kernel parameters and the other dismissing template kernel parameters and the performance of both will be sketched in the results Section. This has been carried out to cope with different degrees of knowledge at compile time on the problem parameters.

It should be mentioned that all the host (__host__) or device (__global__ and __device__) functions involved in the implementation have been designed as template functions so that the same approach can be exploited for single or double precision arithmetics.

It should also be underlined that, in the code snippets that will illustrate the approach, for the sake of clarity, any type of template as well as keywords like const and __restrict__ will be omitted.

3.1. Data Organization in Device Global Memory

Being the algorithm memory bound, one bottleneck of the GPU implementation is represented by the data memory access and organization, which should be set up so that the threads access the global memory in a coalesced way. This memory setting is exploited by the first part of the approach. In order to describe how this point has been here faced, let us assume that K is the number of matrices for which the SVs are required, namely

$$\begin{array}{ccc}\hfill A_1& =& [q_1^{\left(1\right)},q_2^{\left(1\right)},\dots ,q_n^{\left(1\right)}]\hfill \\ \hfill A_2& =& [q_1^{\left(2\right)},q_2^{\left(2\right)},\dots ,q_n^{\left(2\right)}]\hfill \\ \hfill A_3& =& [q_1^{\left(3\right)},q_2^{\left(3\right)},\dots ,q_n^{\left(3\right)}],\hfill \\ \hfill \dots & =& \dots \hfill \\ \hfill A_K& =& [q_1^{\left(K\right)},q_2^{\left(K\right)},\dots ,q_n^{\left(K\right)}]\hfill \end{array}$$

(24)

where $q_j^{\left(i\right)}$, for $i=1,\dots ,K\mathrm{and}j=1,\dots ,n$, represents the jth column of the ith matrix.

Henceforth, we will assume a column-major ordering of each matrix so that the matrix storage becomes compatible with the storage scheme conventionally used in standard libraries, such as cuBLAS, for example. Furthermore, we will assume that, in the first step of the procedure (bidiagonalization), which is the step when all the K matrices must be read from global memory, each GPU thread is assigned to a different matrix. This assumption will be clarified in the next Subsection.

A few words on coalesced memory transactions are now in order. The device code is actually executed in groups of 32 threads, which are called warps. When a warp executes an instruction that contains a global memory access, it automatically acquires the content of consecutive global memory locations into a single transaction. Accordingly, if the threads access the global memory in a consecutive way (coalesced memory access), then the data acquired by the memory transaction are exploited and the bandwidth is efficiently used. If, on the other side, the threads do not access memory in a consecutive way, then there can be unused words of the transaction that are transferred from/to global memory in addition to the words used by the threads, thus reducing the memory throughput. The more unused words are accessed, the more the throughput is reduced, resulting in a very inefficient code, of course.

Figure 4 illustrates a natural, consecutive memory storage of the pool of matrices. With each thread assigned to a different matrix, such a storage scheme does not lead to a coalesced memory access since consecutive threads do not access consecutive regions of memory. Opposite to that, a more performing data organization, which has been used in this paper, consists of an interlaced matrix storage as illustrated in Figure 5. In this case, parallel threads running the same instruction access consecutive locations of global memory, leading to a far higher memory bandwidth.

In other words, Figure 4 demonstrates how naïve back-to-back storage causes each warp to fetch scattered addresses, leading to under-utilized memory transactions. By contrast, Figure 5’s interlaced layout aligns consecutive thread IDs with consecutive addresses, enabling full 128-byte coalesced reads. To better clarify this point, Figure 6a illustrates K input matrices stored one after another as K stacks of bricks, where each brick is one matrix. In the conventional layout, continuous GPU threads uselessly pick up inhomogeneous stacks, and memory transactions fetch many unused words. In the interlaced layout illustrated in Figure 6b, the bricks are redistributed so that each thread takes homogeneous bricks. In other words, threads in a warp access consecutive bricks, all of which are used, yielding coalesced global-memory reads.

According to what the above said, and in order to interface the natural input storage of Figure 4 to that of Figure 5, the __global__ function in Algorithm 2 is employed.

Algorithm 2 Memory reorganization __global__ function.

__global__void $reorganize$ (double *$in$, double *$out$, int                                               K, int $rows$, int $cols$) {       int $tid=threadIdx.x+blockDim.x\ast blockIdx.x$;       if $(tid<K)$           for (int $i=0;i<rows;i$++)                for (int $j=0;j<cols;j$++)                    $out[tid+j\ast K+i\ast K\ast cols]=$                      $in[tid\ast rows\ast cols+j\ast rows+i];$ }

3.2. Thread-Level and Instruction-Level Parallelism

Before describing the parallelization strategies devised for steps 1 through 3, we mention the opportunity of using thread-level parallelism or instruction-level parallelism.

The achievable level of thread parallelism with very tiny matrices, as in the case discussed in the present work, is very low due to the organization of the execution pattern on the GPU. Accordingly, the parallelization strategies for steps 1 through 3 will rely much on the exploitation of instruction-level parallelism.

3.3. Parallelization of the First Step (Bidiagonalization)

Computing the bidiagonalization with the scheme discussed in Section 2.1, namely, using a series of Householder transformations, has two main drawbacks for parallel implementation on GPUs.

The first, most evident drawback is the data dependency between each bidiagonalization step. Indeed, the right lower $(m-i)\times (n-i)$ submatrix of $A_i$ obtained at the ith iteration step of the procedure is the starting point for the calculation of the submatrix $A_{i+1}$ at the $(i+1)$th step.

The second issue concerns the continuous reduction of the matrix size against the step index i, which implies that each iteration operates on a smaller input matrix as compared to the previous one. Accordingly, a continuous reduction of the workload takes place. Continuously reducing the workload is an issue for parallel implementations on GPUs since it is likely to prevent saturating the computational resources of the device.

Furthermore, concerning the matrix update involved in Equation (10), it should be mentioned that the popular approach in [46] partitions the input matrix into square blocks. Assuming that L is the dimension of each block, the matrix update occurs only after L rows and L columns have been bidiagonalized. In this way, the matrix update is deferred, avoiding the update after each row-column elimination which involves many read/write memory operations. Unfortunately, this approach has been devised for large matrices and is not suitable for matrices as small as those dealt with in this paper.

Following the above considerations and the one in the foregoing subsection, in our approach, each thread is in charge of computing the whole bidiagonalization for a single matrix. Therefore, if K is the size of the pool of matrices, the kernel launch configuration is organized as a one dimensional grid of thread blocks with a fixed number of threads per block equal to $BLOCKSIZE$ and a number of blocks equal to the smallest integer not smaller than $K/BLOCKSIZE$.

The bidiagonalization step has been implemented by the kernel function reported in Algorithm 3. computeLeftHouse and computeRightHouse are __device__ calculating left and right Householder’s vectors, respectively, according to Section 2.1.

Algorithm 3 Bidiagonalization __global__ function.

__global__void $bidiag$ (double *$in$, int $maxit$, int K, int $rows$, int $cols$) {      int $tid=threadIdx.x+blockDim.x\ast blockIdx.x$;        double $leftHouse\left[maxrows\right]$, $rightHouse\left[maxcols\right]$, $h\left[maxcols\right]$, $wi\left[maxrows\right]$, $xi\left[maxcols\right]$, $zi\left[maxcols\right]$;      double $betau=betav=0.$, $sum$;      for (int $count=0;count<maxit;count$++) {          if ($count<cols-2$) {                  if ($tid<K$) {                    $computeLeftHouse$($in$, $leftHouse$, $betau$, $tid$, $count$, K, $rows$, $cols$);                    for (int $j=0;j<cols;j$++) {                      $sum=0.$;                      for (int $i=0;i<rows;i$++) $sum+=-betau\ast leftHouse\left[i\right]\ast in[tid+i\ast K\ast cols+j\ast K]$;                      $h\left[j\right]=sum+in[tid+j\ast K+count\ast K\ast cols]$;                      $xi\left[j\right]=-sum$; }                    $computeRightHouse$(h, $rightHouse$, $betav$, $tid$, $count$, K, $rows$, $cols$);                    for (int $i=0;i<rows;i$++) {                      $sum=0.$;                      for (int $j=0;j<cols;j$++) $sum+=betav\ast rightHouse\left[j\right]\ast in[tid+i\ast K\ast cols+j\ast K]$;                      $wi\left[i\right]=sum$; }                    $sum=0.$;                    for (int $i=0;i<cols;i$++) $sum+=xi\left[i\right]\ast rightHouse\left[i\right]$;                    for (int $i=0;i<cols;i$++) $zi\left[i\right]=xi\left[i\right]-betav\ast sum\ast rightHouse\left[i\right]$;                    for (int $r=0;r<rows;r$++)                      for (int $c=0;c<cols;c$++) $in[tid+r\ast K\ast cols+c\ast K]=-leftHouse\left[r\right]\ast zi\left[c\right]$; }}                  if ($count>=cols-2$) {                    if ($tid<K$) {                    $computeLeftHouse$($in$, $leftHouse$, $betau$, $tid$, $count$, K, $rows$, $cols$);                    for (int $j=0;j<cols;j$++) {                      double $sum=0.$;                      for (int $i=0;i<rows;i$++) $sum+=leftHouse\left[i\right]\ast in[tid+i\ast K\ast cols+j\ast K]$;                      $xi\left[j\right]=sum$; }                    for (int $r=0;r<rows;r$++)                      for (int $c=0;c<cols;c$++) $in[tid+r\ast K\ast cols+c\ast K]=-betau\ast xi\left[c\right]\ast leftHouse\left[r\right]+in[tid+r\ast K\ast cols+c\ast K];$ } } } }

3.4. Parallelizations of the Second (Tridiagonalization) and Third (Root Finding) Steps

Step 2, namely tridiagonalization, is trivially implemented as a __device__ function in CUDA, so that this point will be here skipped.

Following the considerations in Section 2.3 about the embarrassingly level of parallelism of the third step, the bisection algorithm is executed on a 2D grid of threads, with the threads along the x dimension assigned to a different matrix and the threads along the y dimension in charge of calculating a single eigenvalue of each matrix.

The bisection approach is implemented by the kernel function reported in Algorithm 4. The functions chcksign, pivot, tridiag, and initInterv appearing there are __device__ functions. In particular, chcksign and pivot implement the Count function (see Appendix A), tridiag calculates the elements of the tridiagonal matrix, and initInterv calculates the initial search interval for bisection.

Algorithm 4 Bisection __global__ function.

__global__void $bisection$ (double * d, double * b, double * $dd$, double *$bb$, double *$SV$, int K, int $cols$, double $tol$) {         int $tidx=threadIdx.x+blockDim.x\ast blockIdx.x$;         int $tidy=threadIdx.y$;         double $alpha$, $beta$, $pivot$;           __shared__double   $alph{a}_{shared}\left[blockSizeSturm\right]$, $bet{a}_{shared}\left[blockSizeSturm\right]$, $pivmi{n}_{shared}\left[blockSizeSturm\right]$;         if ($tidx<K$ && $tidy==0$) {             $tridiag$($d+tidx$, $b+tidx$, $dd+tidx$, $bb+tidx$, K, $cols$);             $pivot$($bb+tidx$, $\&pivmi{n}_{shared}[threadIdx.x]$, K, $cols$);             $initInterv$($dd+tidx$, $bb+tidx$, $alph{a}_{shared}[threadIdx.x]$, $bet{a}_{shared}[threadIdx.x]$, K, $cols$); }         __syncthreads();         if ($tidx<K$ && $tidy<cols$) {           int $eigIndex=tidy\%cols$;             int $numChanges=0$;             double $c=0$;             $alpha=alph{a}_{shared}[threadIdx.x]$;             $beta=bet{a}_{shared}[threadIdx.x]$;             $pivot=pivmi{n}_{shared}[threadIdx.x]$;             double $dist=fabs(beta-alpha)$;             double $s=fabs\left(beta\right)+fabs\left(alpha\right)$;             while ($dist>tol\ast s$) {                 $c=(alpha+beta)/2.$;                 $chcksign$($dd+tidx$, $bb+tidx$, $pivot$, c, $numChanges$, K, $cols$);                 if $(numChanges>(cols-(eigIndex+1)\left)\right)$$beta=c$;                 else $alpha=c$;                 $dist=fabs(beta-alpha)$;                 $s=fabs\left(beta\right)+fabs\left(alpha\right)$; }             $SV[eigIndex+tidx\ast cols]=sqrt\left(fabs\right(c\left)\right)$; } }

3.5. Multi-GPU Processing

The multi-GPU processing is implemented by distributing the overall number $K_T$ of matrices to be processed on the available $N_{GPU}$ GPUs. In this way, each GPU operates on a number of $K=K_T/N_{GPU}$ matrices. The data transfers towards the GPUs is performed by asynchronous copies using cudaMemcpyAsync to make them occur concurrently.

3.6. Computational Complexity and Memory Requirements Per-Thread

3.6.1. Bidiagonalization

The cost of the bidiagonalization via Householder reflections step per matrix, namely, per thread, is $\simeq (m·n^2)$ [47]. The GPU global memory requirement per-thread in bytes is $\left(m\right(n+2)-1)\mu \left(T\right)$, where $\mu \left(T\right)=4$ for single precision or $\mu \left(T\right)=8$ for double precision. The GPU local per thread register/stack memory requirement is $(3m+4n)\mu \left(T\right)$.

3.6.2. Tridiagonalization

The cost of the tridiagonalization step is $\simeq (m+n)$. The GPU global memory requirement per thread in bytes is $2(m+n)\mu \left(T\right)$.

3.6.3. Root-Finding by Sturm Bisection

The cost of the root-finding by bisection with Sturm sequences per matrix is $\mathcal{O}(nlog(W/\delta \left)\right)$, where W is the Gershgorin-interval width and $\delta$ the tolerance. The GPU global memory requirement per thread in bytes is $(m+2(n+1\left)\right)\mu \left(T\right)$.

3.7. Comparison with the Approaches in [14,22,23]

The approach in [22] has a computational complexity of $\simeq \left(Sm{n}^2\right)$, where S is the number of sweeps which can be, in the worst case, $\simeq log\left(n\right)$. Opposite to that, the computational complexity of the scheme in [23] is $\simeq \left(m{n}^2\right)$. Finally, for square matrices, the approach in [14] is $\simeq \left(n^3\right)$.

4. Numerical Results

In this section, we present some of the results of an extensive numerical analysis aimed at assessing the performance of the proposed approach. Although the CUDA code has been implemented in both single- and double-precision arithmetics, double precision is mainly considered here, with most of the trends for single precision being similar.

Performance is assessed against a different size $K_T$ of the batch to be processed. In particular, for all the examples below, $K_T=2^n$, $n=4,\dots ,20$. In this way, either the case of a small and of a large number of matrices can be dealt with. Furthermore, we consider the case of either square or rectangular matrices. In the rectangular case, we set $m=n-2$ to have again the possibility to evaluate the performance for an increasing matrix size. Below, we separate the analysis between the cases of single-GPU and multi-GPU processing.

We finally underline that the GPU timings below include the bidirectional host-device memory transfers. The plots are all reported in log-log scale.

The assessment has been performed on different single- and multi-GPU architectures. We will firstly consider the case of a g4dn.12xlarge AWS instance (Amazon SageMaker) comprising 48 vCPUs, CPUs of the Intel Xeon Family, 192 GB of memory, 4 NVIDIA T4 Tensor Core GPUs of compute capability $7.5$, and 64 GB of global memory. The employed CUDA version has been $12.2$ and the code has been compiled with the $O3$ optimization option.

4.1. Accuracy, Runtime, and Stability

Before assessing the computational performance of the technique, we analyze its accuracy and stability. In particular:

• we quantify the effect of the tolerance parameter $\delta$ on both accuracy and runtime;
• we assess the robustness of the method on ill-conditioned matrices;
• we assess the robustness of the method against noisy matrices.

4.1.1. Percentage Root Mean Square Error

Let A be an $m\times n$ test matrix, $\sigma =\left\{{\sigma }_i\right\}$ the vector of the singular values computed by our approach, and $\widehat{\sigma }=\left\{{\widehat{\sigma }}_i\right\}$ the reference singular values. The percentage relative root mean square SV error (PRMSE) is computed as

$$E_{\mathrm{RMSE}}=100·\sqrt{{\displaystyle \frac{\parallel \sigma -\widehat{\sigma }{\parallel }_2}{\parallel \widehat{\sigma }{\parallel }_2}}},$$

(25)

where ${\parallel ·\parallel }_2$ denotes the standard quadratic norm.

4.1.2. Tolerance Against Accuracy and Runtime

Figure 7 shows an accuracy/runtime map for single- and double-precision arithmetic for random Gaussian matrices having $3\times 3$, $4\times 4$, $5\times 5$, and $6\times 6$ sizes and for different tolerance parameters $\delta$. In particular, $\delta ={10}^{-6}$ and $\delta ={10}^{-7}$ have been dealt with for single-precision, while $\delta ={10}^{-12}$ and $\delta ={10}^{-13}$ have been considered for double-precision. As can be seen, the resulting RMSE agrees with the typical tolerances of the two kinds of arithmetic, and, of course, more accuracy is paid with a larger computation time.

4.1.3. Stability Against the Condition Number

We have considered $4\times 4$ Gaussian random matrices with different condition numbers and double-precision arithmetics, having fixed $\delta ={10}^{-13}$. In particular, the matrices have condition numbers $D={10}^{\kappa }$, with $\kappa \in \{2,4,6,8,10\}$, and have been defined by generating random singular vector matrices and singular values $d_i={10}^{\kappa (i-1)/(m-1)}$, $i=0,\dots ,m-1$.

Table 1. Stability against the condition number.

D	PRMSE
${10}^2$	$5.3\times {10}^{-12}$
${10}^4$	$1.74\times {10}^{-10}$
${10}^6$	$3.86\times {10}^{-9}$
${10}^8$	$5.07\times {10}^{-8}$
${10}^{10}$	$1.42\times {10}^{-6}$

summarizes the results of the analysis and shows how the PRMSE keeps tolerable even for very highly ill-conditioned matrices.

4.1.4. Noise Perturbation

Finally, to probe sensitivity to noise, we have considered noisy, $4\times 4$ matrices using $A_{\mathrm{noisy}}=A+ϵ{\parallel A\parallel }_{2}G$, with G Gaussian and $ϵ\in \{{10}^{-8},{10}^{-7},{10}^{-6},{10}^{-5}\}$, again with double-precision arithmetic and $\delta ={10}^{-13}$.

Table 2. Stability against the noise.

$\mathit{ϵ}$	PRMSE
${10}^{-8}$	$2.42\times {10}^{-12}$
${10}^{-7}$	$2.91\times {10}^{-12}$
${10}^{-6}$	$2.54\times {10}^{-12}$
${10}^{-5}$	$2.72\times {10}^{-12}$

resumes the stability of the algorithm against the noise.

4.2. Performance of the Approach for the Single GPU Case

Figure 8 illustrates the processing time for the proposed approach when only 1 GPU is employed. The processing time is reported against the number of matrices and parameterized on the size of the involved matrices. The template version of the code has been run with a tolerance $\delta$ chosen equal to ${10}^{-7}$. It can be seen that, unless the batch saturates the GPU resources, which occurs for $K_T\simeq 1024$, the processing time does not significantly change with the batch size.

Figure 9, on the other side, illustrates the separate timings of the different algorithm stages. It highlights how the most computationally expensive stage is represented by the tridiagonalization and bisection one.

We have also confronted our scheme with a fully CPU one based on the use of the parallel, CPU multicore Intel MKL library [10]. Comparing the proposed approach with a multicore CPU implementation has a rationale in the fact that multicore CPUs are often perceived as the best computing platform for such kinds of computational problems, namely, computations on very large numbers of small items. However, also in this case, our algorithm has outperformed the compared one.

We now want to illustrate the speedup obtained by the above-discussed template implementation of the code. More in detail, as can be seen from Figure 10, the template version enables gaining about $10\%$ of the overall processing time against the non-template version of the approach.

We finally want to show, in Figure 11, how the algorithm performs for non-square matrices. As it can be seen, there is no significant performance loss due to the rectangularity of the matrices. Also, we have compared the computing time of the case when $\delta ={10}^{-7}$ against the case $\delta ={10}^{-9}$. A performance loss of about $10\%$ has been observed, which is definitely tolerable in applications requiring more accuracy: despite the requested accuracy improvement of two orders of magnitude, the case $\delta ={10}^{-7}$ is not much faster than the case $\delta ={10}^{-9}$.

Finally,

Table 3. Accuracy results.

Size	MKL	Double $\mathit{\delta }={10}^{-7}$	Double $\mathit{\delta }={10}^{-9}$	Float $\mathit{\delta }={10}^{-4}$	Size	Double $\mathit{\delta }={10}^{-7}$
$2\times 2$	$3.2\times {10}^{-14}$	$1.7\times {10}^{-13}$	$1.65\times {10}^{-17}$	$1.89\times {10}^{-7}$	$4\times 2$	$1.54\times {10}^{-13}$
$3\times 3$	$9.18\times {10}^{-14}$	$1.9\times {10}^{-13}$	$1.45\times {10}^{-17}$	$1.99\times {10}^{-7}$	$5\times 3$	$1.8\times {10}^{-13}$
$4\times 4$	$4.88\times {10}^{-14}$	$2.12\times {10}^{-13}$	$1.37\times {10}^{-17}$	$2.13\times {10}^{-7}$	$6\times 4$	$2.09\times {10}^{-13}$
$5\times 5$	$3.87\times {10}^{-14}$	$2.22\times {10}^{-13}$	$1.45\times {10}^{-17}$	$2.30\times {10}^{-7}$	$7\times 5$	$2.14\times {10}^{-13}$
$6\times 6$	$4.38\times {10}^{-14}$	$2.22\times {10}^{-13}$	$1.47\times {10}^{-17}$	$2.41\times {10}^{-7}$	$8\times 6$	$2.21\times {10}^{-13}$
$7\times 7$	$7.89\times {10}^{-14}$	$2.35\times {10}^{-13}$	$1.52\times {10}^{-17}$	$2.47\times {10}^{-7}$	$9\times 7$	$2.23\times {10}^{-13}$
$8\times 8$	$1.21\times {10}^{-13}$	$2.22\times {10}^{-13}$	$1.53\times {10}^{-17}$	$2.33\times {10}^{-7}$	$9\times 7$	$2.27\times {10}^{-13}$

reports the root-mean-square (RMS) errors of the singular values computed by our GPU-based method versus MATLAB’s state-of-the-art svd reference, for square matrices with sizes from $2\times 2$ to $8\times 8$, and for two tolerance levels for double precision arithmetic, namely, $\delta ={10}^{-7}$ and $\delta ={10}^{-9}$, and for a single tolerance level for single-precision arithmetic, namely, $\delta ={10}^{-4}$. Furthermore, Table 3 also reports the RMS errors for some rectangular matrices for $\delta ={10}^{-7}$ and double precision arithmetic only. For each size in the square matrices case, we also include the RMS errors obtained with Intel MKL’s desvg routine under its default settings. As shown, when $\delta ={10}^{-7}$, our approach already achieves $<{10}^{-12}$ RMS errors, closely matching MATLAB’s results. Tightening the tolerance to $\delta ={10}^{-9}$ reduces the RMS error further—dropping below ${10}^{-16}$ and outperforming MKL. This demonstrates that our bisection-based solver not only meets the accuracy of established CPU routines but also allows explicit control of precision via the $\delta$ parameter without significant loss of robustness.

4.3. Performance of the Approach in Comparison to cuSOLVER’s Gesvd

The performance of the approach is further compared to cuSOLVER’s gesvd routine on four different GPUs, namely, A10, A6000, RTX6000, and T4, for the case of double precision arithmetic. In particular, Figure 12 and Figure 13 illustrate the speedup experienced when using the proposed numerical scheme for $4\times 4$ and $10\times 10$ matrices in the case of double precision arithmetic. The advantage of our technique is highlighted by a speedup reaching 20 for batches of $4\times 4$ matrices and $3.5$ for batches of $10\times 10$ matrices. The PRMSE in our case has been around $4\times {10}^{-12}$, while that for cuSOLVER’s gesvd routine has been also lower settling down around $3\times {10}^{-7}$. We notice that cuSOLVER’s gesvd forms the basis for several, publicly available routines as current versions of JAX, PyTorch, and CuPy.

4.4. Performance of the Approach for the Multi-GPU Case

Finally, in order to assess the scaling capabilities of our scheme in the multi-GPU case, Figure 14 and Figure 15 show the speedup achieved against the single-GPU case when 2 or 4 GPUs are employed, respectively. As can be seen, once the GPU resources have been saturated, scalings of about 2 and 4 are observed for the two cases.

To further confirm the scaling features of the approach, we have tested it against a dual-GPU system based on NVIDIA A6000 cards. More in detail, we have considered a batch of ${10}^{23}$ matrices. The case of $4\times 4$ matrices run in $0.32$ s on a single GPU and in $0.16$ s on two GPUs, while the case of $10\times 10$ matrices took $2.25$ s to run on a single GPU and $1.13$ s to complete the processing on two GPUs, thus confirming the good scaling features of our approach.

5. Conclusions and Future Developments

A fast and robust approach to evaluate the singular values of small matrices on Graphics Processing Units (GPUs) was presented, exploiting interlaced storage, bidiagonalization via Householder transformations, symmetric tridiagonalization, and Sturm-sequence–based bisection. The algorithm was implemented in CUDA and tested on several single- and multi-GPU architectures. For large matrix batches, the speed-up against cuSOLVER’s gesvd has been around 20 for $4\times 4$ matrices. Furthermore, near-linear scaling across the GPUs ($\simeq 4\times$ acceleration) has been yielded, while maintaining an RMS error below $2.3\times {10}^{-7}$ in single precision for relative tolerance $\delta ={10}^{-4}$ and below $2.3\times {10}^{-13}$ in double precision for $\delta ={10}^{-7}$. By tightening the tolerance to $\delta ={10}^{-9}$, RMS errors dropped below $2\times {10}^{-13}$, thus matching the accuracy of MATLAB’s state-of-the-art svd routine and demonstrating that, with proper choice of $\delta$, the proposed method can achieve precision equivalent to MATLAB’s svd. Furthermore, tightening the tolerance from $\delta ={10}^{-7}$ to $\delta ={10}^{-9}$ increased the total runtime by only about 10%. It was also verified that the performance in terms of execution time for rectangular matrices did not differ significantly as compared to the square matrices case. Future developments include extension to complex matrices and application to large-scale singular value optimization problems. Steps of the algorithm requiring the computation of matrix multiplications can also be further improved using approaches specifically tailored to small matrices as in [14].