CS355 CUDA programming project

Preparation
- Execute the following command ONCE to setup the CUDA environment:
  Logout and then log back in to update your shell environment variables.
- Execute the following command to copy the necessary files into your project directory:

Assignment

In this assignment, you will write a CUDA kernel to compute the inverse of a NxN matrix

To help you do this project, I have provided:

Extensive review of material from Math221 (Linear Algebra) to refresh your knowledge on matrix inversion
The review material is in the 2 sections in this handout below - click here to go to the review section
A CPU function invert( ) in the C program file invert.c (made available in the preparation step above) that contain the Math221 (Linear Algebra) algorithm to computes the inverse of a matrix

What you need to do for the CUDA project:

Edit the file cuda-inverse.cu and write the inverse( ) kernel:

/* ======================================================================= TODO: write a CUDA kernel to compute the inverse of matrix A inverse(A, C, N): input: A = an NxN matrix that you need to find the inverse When the inverse( ) function completes, A will contains the identity matrix C = initially contains the identity matrix When the inverse( ) function completes, C will contains the inverse of A N = #row (and # columns) in A and C ======================================================================== */ __global__ void inverse( float *A, float *C, int N) { // You need to write this function for the CUDA oroject // I have given some hints in the file to help you write // this CUDA kernel }

How to compile and run the project

Important:

You must login to lab1X (X is one of a, b, c, d, e, f, g, h) to do the work:

From lab0z or some lab machine: ssh -X yourNetID@lab1a or: ssh -X yourNetID@lab1b etc....

To compile (on lab1X):

/usr/local/cuda/bin/nvcc -o cuda-inverse cuda-inverse.cu

To run your project (on lab1X):

cuda-inverse N or cuda-inverse N manual // Manual entry of NxN matrix

Important:

The running time for a 1000x1000 matrix of your CUDA program must be less than the CPU version (that I provide) !!!
(If your running time is longer, than you have probably not written a parallel program -- serial code is easy to write; you must learn to write parallel code in this project)

Sample outputs

100x100 matrix:

Inverting a 100x100 matrix: cs355@lab1X (3273)> inverse 100 Elasped time = 16794 micro secs Max err = 0.000208 Inverse is correct ! cs355@lab1X (3274)> cuda-inverse 100 Elasped time = 6477 micro secs N > 5; skip printing... Max err = 0.000133 Inverse is correct !

1000x1000 matrix:

Inverting a 1000x1000 matrix: cs355@lab1X (3275)> inverse 1000 Elasped time = 7262190 micro secs Max err = 0.055180 Inverse is correct ! cs355@lab1X (3276)> cuda-inverse 1000 Elasped time = 4011005 micro secs N > 5; skip printing... Max err = 0.018559 Inverse is correct !

Inverse of a matrix (from Math221 - Linear Algebra)

Inverse of a matrix (Linear Algebra):

Given matrix A A^-1 is the inverse of matrix A if and only if: A*A^-1 = I where I = the identity matrix

Example:

+- -+ | 7 2 1 | A = | 0 3 -1 | | -3 4 -2 | +- -+
The inverse of A is: +- -+ | -2 8 -5 | B = | 3 -11 7 | | 9 -34 21 | +- -+ Because: +- -+ +- -+ | 7 2 1 | | -2 8 -5 | A*B = | 0 3 -1 | * | 3 -11 7 | | -3 4 -2 | | 9 -34 21 | +- -+ +- -+ +- -+ | -7*2+2*3+1*9 7*8-2*11-1*34 -7*5+2*7+1*21 | = | -0*2+3*3-1*9 0*8-3*11+1*34 -0*5+3*7-1*21 | | 3*2+4*3-2*9 -3*8-4*11+2*34 3*5+4*7-2*21 | +- -+ +- -+ | 1 0 0 | = | 0 1 0 | <--- identify matrix | 0 0 1 | +- -+

How to find the inverse of a matrix

In Math221, you have learned the following Gauss-Jordan algorithm to find the inverse of a NxN matrix:

Gauss-Jordan Method to find the inverse of:

+- -+ | 7 2 1 | A = | 0 3 -1 | | -3 4 -2 | +- -+

We first setup the matrix A,I:

+- -+ +- -+ | 7 2 1 | | 1 0 0 | A , I = | 0 3 -1 | | 0 1 0 | | -3 4 -2 | | 0 0 1 | +- -+ +- -+ Simplified representation: 7.000 2.000 1.000 | 1.000 0.000 0.000 0.000 3.000 -1.000 | 0.000 1.000 0.000 -3.000 4.000 -2.000 | 0.000 0.000 1.000

Row 0:

Multiply row₀ by 1/A[0][0]; // Normalize row 0 of matrix; A[0][0] = 7 --> divide row 0 by 7 Result: 1.000 0.286 0.143 | 0.143 0.000 0.000 <--- Updated row 0 0.000 3.000 -1.000 | 0.000 1.000 0.000 -3.000 4.000 -2.000 | 0.000 0.000 1.000
// Use row 0 to create a column of 0 values in column 0 What you will be doing are: add 0x row 0 to row 1 // "Reduce" row 1 add +3x row 0 to row 2 // "Reduce" row 2 1.000 0.286 0.143 | 0.143 0.000 0.000---+ 0.000 3.000 -1.000 | 0.000 1.000 0.000<--+ -0x row 0 -3.000 4.000 -2.000 | 0.000 0.000 1.000 After reducing row 1: 1.000 0.286 0.143 | 0.143 0.000 0.000---+ 0.000 3.000 -1.000 | 0.000 1.000 0.000 | -3.000 4.000 -2.000 | 0.000 0.000 1.000<--+ +3x row 0 After reducing row 2: 1.000 0.286 0.143 | 0.143 0.000 0.000 0.000 3.000 -1.000 | 0.000 1.000 0.000 0.000 4.857 -1.571 | 0.429 0.000 1.000

Row 1:

Multiply row₁ by 1/A[1][1]; // Normalize row 1 of matrix; A[1][1] = 3 --> divide row 1 by 3 Result: 1.000 0.286 0.143 | 0.143 0.000 0.000 0.000 1.000 -0.333 | 0.000 0.333 0.000 <--- Updated row 1 0.000 4.857 -1.571 | 0.429 0.000 1.000
// Use row 1 to create a column of 0 values in column 1 What you will be doing are: add -0.286x row 1 to row 0 // "Reduce" row 0 add -4.857x row 1 to row 2 // "Reduce" row 2 1.000 0.286 0.143 | 0.143 0.000 0.000<--+ -0.286x row 1 0.000 1.000 -0.333 | 0.000 0.333 0.000---+ 0.000 4.857 -1.571 | 0.429 0.000 1.000 After reducing row 0: 1.000 0.000 0.238 | 0.143 -0.095 0.000 0.000 1.000 -0.333 | 0.000 0.333 0.000---+ 0.000 4.857 -1.571 | 0.429 0.000 1.000<--+ -4.857x row 1 After reducing row 2: 1.000 0.000 0.238 | 0.143 -0.095 0.000 0.000 1.000 -0.333 | 0.000 0.333 0.000 0.000 0.000 0.048 | 0.429 -1.619 1.000

Row 2:

Multiply row₂ by 1/A[2][2]; // Normalize row 2 of matrix; A[2][2] = 0.048 --> Divide row 2 by 0.048 Result: 1.000 0.000 0.238 | 0.143 -0.095 0.000 0.000 1.000 -0.333 | 0.000 0.333 0.000 0.000 0.000 1.000 | 9.000 -34.000 21.000 <--- Updated row 2
// Use row 2 to create a column of 0 values in column 2 What you will be doing are: add -0.238x row 2 to row 0 // "Reduce" row 0 add +0.333x row 2 to row 1 // "Reduce" row 1 1.000 0.000 0.238 | 0.143 -0.095 0.000<--+ -0.238x row 2 0.000 1.000 -0.333 | 0.000 0.333 0.000 | 0.000 0.000 1.000 | 9.000 -34.000 21.000---+ After reducing row 0: 1.000 0.000 0.000 | -2.000 8.000 -5.000 0.000 1.000 -0.333 | 0.000 0.333 0.000<--+ +0.333x row 2 0.000 0.000 1.000 | 9.000 -34.000 21.000---+ After reducing row 1: 1.000 0.000 0.000 | -2.000 8.000 -5.000 0.000 1.000 0.000 | 3.000 -11.000 7.000 0.000 0.000 1.000 | 9.000 -34.000 21.000 ^^^^^^^^^^^^^^^^^^^^^ ^^^^^^^^^^^^^^^^^^^^^ DONE: Identity matrix Inverse matrix

Result: inverse matrix is

-2.00 8.00 -5.00 3.00 -11.00 7.00 9.00 -34.00 21.00

A CPU version of the matrix inversion algorithm

In the file inverse.c, I have written an inverse( ) function that implements the Linear Algebra matrix inversion algorithm:

void inverse( float *A, float *C, int N) { /* ========================================================== R = a variable that goes through each row of the matrix ========================================================== */ for (int R = 0; R < N; R++) { float factor = A[R*N + R]; // Use A[R][R] as multiply factor for (int j = 0; j < N; j++) // Normalize row R with factor { A[R*N+j] = A[R*N+j]/factor; C[R*N+j] = C[R*N+j]/factor; } /* ========================================================= Make a column of 0 values in column R using the row "R" ========================================================= */ for (int i = 0; i < N; i++) { if (i == R) { // Do nothing to row "R" } else { float f = A[i*N+R]; // Multiply factor /* ------------------------------------- Add -f*row(R) to row(i) ------------------------------------- */ for (int j = 0; j < N; j++) { A[i*N+j] = A[i*N+j] - f*A[R*N+j]; C[i*N+j] = C[i*N+j] - f*C[R*N+j]; } } } } }

You can run this program as follows:

Compile: gcc -o inverse inverse.c Run: inverse N or: inverse N manual // Manual entry of a NxN matrix

Turn in

Turn in your CUDA project using the following command on a lab machine before the deadline:

cd ~/cs355/CUDA /home/cs355001/turnin cuda-inverse.cu cuda

Extension request
- If you want to request an extension for this project, execute the follow command on a lab machine: