# Calculating the Moore-Penrose Generalized Inverse on Massively Parallel Systems
## 3 Massively Parallel Computing and CUDA  
GPUs are traditionally used for the fast generation and manipulation of images by splitting the massive computation task into tiny operations on a small volume of data performed in parallel threads. Usually, each thread is related to a single pixel. This technology was adjusted for parallel computing on many cores. Using GPUs for general-purpose computing was made accessible in 2006 when NVidia released the first version of the CUDA platform [39] with indirect access to the instructions and memory for parallel computation on GPUs. The code of a CUDA program contains the host code that is executed on a CPU and the device code that is executed on a GPU. The host and the device have distinct memories and address spaces, and the device memory is managed by built-in functions. In most cases, we can identify three main steps in the execution flow of a CUDA program: copying the memory from the CPU to GPU, the function calls and the execution of functions that are executed on the GPU, and copying the results from the GPU to CPU [39]. The procedures and functions that are executed on the GPU in parallel are called kernels (keyword_global_ in their declaration). The most important factor in a parallel version of the algorithm is the execution of the most expensive calculations using the kernel functions and avoiding heavy information traffic between the GPU and the host device. The kernel body, using the CUDA programming model, is executed in parallel with many threads grouped in blocks, which are further grouped in a grid. The maximum number of threads one block can contain is 1024. The CUDA allows blocks and grids to be organized in one, two, or three dimensions. Every thread has a unique identifier within the block, which can be accessed through the built-in _threadIdx_, which contains \(x,y,\) and \(z\) attributes. Similarly, each thread can access the _blockIdx_ variable. The dimensions of the block and grid can be accessed through the built-in variables _blockDim_ and _gridDim_. We use the following formulas to identify a thread within a block and a block within a grid:  
\[uniqueThreadIdWithinBlock=x_{t}+y_{t}B_{x}+z_{t}B_{x}B_{y}, \tag{4}\]  
\[uniqueBlockId=x_{b}+y_{b}G_{x}+z_{b}G_{x}G_{y}, \tag{5}\]  
where (\(B_{x}\), \(B_{y}\), \(B_{z}\)) and (\(G_{x}\), \(G_{y}\), \(G_{z}\)) are the dimensions of the block and grid, respectively, and \(x_{t},y_{t}\), and \(z_{t}\) and \(x_{b}\), \(y_{b}\), and \(z_{b}\) represent the _threadIdx.x_, _threadIdx.y_, _threadIdx.z_, _blockIdx.x_, _blockIdx.y_, and _blockIdx.z_, respectively. We can easily calculate the unique global thread ID as:  
\[uniqueGlobalThreadId=N_{b}b_{id}+t_{id}, \tag{6}\]  
where \(b_{id}\) and \(t_{id}\) represent the unique thread ID within the block and block ID, respectively. The GPU contains various types of programmable and nonprogrammable memories. L1 cache, constant cache, L2 cache, and read-only cache are types of nonprogrammable GPU memory. The programmable types of GPU memory are global memory, shared memory, and registers [39]. Global memory is the largest and slowest. Optimizing access to global memory is crucial to improving the program's performance. Shared memory (programmable cache) is a fast memory shared by all the threads within the same block and is visible to all these threads within the same block. It is divided into 32 equally sized banks. Different memory banks can be accessed simultaneously, so it is important to organize the data so that different threads access different shared memory banks. Its address space is one-dimensional, and assuming that we store 4-byte values, the \(k\)th bank contains elements such that \(i\bmod 32=k\), where \(i\) is the index of the data in shared memory. The optimal use of shared memory is crucial for overall performance. In general, the bank index is calculated as  
\[bankIndex=(byteAddress/4)\bmod 32. \tag{7}\]  
The CUDA offers a few types and levels of synchronization.