This repository contains some sample codes helpful in learning how to program with OpenMP and CUDA. Here we are going to go through each directory and explain the purpose of each file. First we are focusing on parallelization using OpenMP. Then we parallelize large calculations using CUDA and GPU. In the end, a sample project for calculating the determinant of large matrices is developed using both OpenMP and CUDA.

In the OpenMP section, there is a sample code in parallel_for_loop.cpp, which, as the name suggests, is a simple for-loop parallelization.

In the matrix_add.cpp code, we have three 2D matrices, A, B, and C, where we want to calculate C = A + B. We do this in two ways: i) row-wise parallelization using a single parallel for-loop and ii) parallelized nested for-loops using the collapse() argument. The elements of A and B matrices are initialized randomly between 0 and 100. The schedule({static, dynamic}, chunk_size) specifies how the thread allocation should be. If the schedule is static, it statically allocates a chunk_size number of successive iterations to a single thread. In the case of dynamic allocation, each thread dynamically selects some number of chunks of iterations and runs them.

In the Mandelbrot section, we have the parallelization of the Mandelbrot set generation, which generates Mandelbrot sets. The serial.cpp file contains the serial code for generating this set. Here we are trying to parallelize this code using OpenMP. The first step is to parallelize the main for-loop in the For.cpp. The second technique is Multi-Tasking which generates multiple tasks to do the process. One major problem that reduces the program speed is writing to file. We can use the Pipeline technique (as we did in Pipeline.cpp) to parallelize both: i) set generation and ii) writing to file.

In the PrefixSum section, we are parallelizing the Prefix sum problem. The serial.cpp contains the serial code for this problem. The naive parallelization of the program creates the challenge of race condition that can reduce the efficiency of parallelism. You can see the first parallelization in parallel_simple.cpp . Another approach to parallelizing this problem is the Hillis and Steele algorithm. Although this algorithm might be slower than the serial mode, it is faster when programmed and run on GPUs.

We start the CUDA section with a test program generated by Visual Studio. The purpose of this program in VS is to ensure that CUDA works. This program in under the VectorAdd directory where we brought the serial code in serial.cpp, the parallelized code using OpenMP in parallel_omp.cpp, and finally the parallel code on GPU in parallel_cuda.cu.

In the OpenMP section, we had a matrix_add.cpp program in which the elements of two matrices A and B are summed and stored in a result matrix C. We have the same program in the CUDA section matrix_add.cu, where we do the same process, except that the matrices are very large, and the parallelization is being done on GPU.

Under the MatrixMult section, we parallelized the matrix multiplication operation. The serial program is in serial.cpp. First, we have a naive parallelization in single_block.cu where we use only a single block of threads in GPU to multiply two matrices. The number of threads in a single block is limited to 1024. Therefore, for matrices with sizes larger than 1024, we should either use tiling method or use multiple blocks. It is not efficient to use a single block of threads while we have multiple free blocks. Therefore, multiple_blocks.cu program does this operation using multiple blocks (without tiling) . We combined the tiling method with multiple blocks in multiple_blocks_tiling.cu. Also, an example of the matrix multiplication operation using the cuBLAS library is given in matmul_cuBLAS.cu.

Another section dedicates to streaming multiple kernels and the concept of pageable and pinned memory. For information about pageable and pinned memory, please visit the following useful link. Also, for streaming, see link. We use the Histogram problem as the example for this section. In the Histogram problem, the goal is to count the number of different elements in a large array. In the pageable.cu we used the pageable memory allocated using the malloc() function. In the pinned.cu, the memory is allocated using cudaMallocHost() and is page-locked, making it faster to access the host memory. In order to use streaming and concurrent kernels, the memory must be pinned. In the streaming.cu we ran multiple kernels in parallel using page-locked memory.

The Determinant directory contains the serial and parallel code for calculating the determinant of large matrices. The method we used is LU decomposition, a simple and fast method for determinant calculation. Sample matrices for testing the code is given in data_in.zip. You should extract this file before running the code in the same directory. The results will be stored in a data_out directory with file names similar to input files except that an outputs_ will be prepended to the name of the input file.

For the parallelization part, reading from and writing to files are done in parallel using OpenMP, while the determinant calculation gets done on GPU.