Goal:

Design a program that computes square matrix multiplication on GPU using CUDA. Write the code in C. In particular, your implementation should obey the following requirements:

1. The program must be general enough to handle matrix sizes beyond the GPU capacity.

2. The GPU capacity should not be hardcoded, but should be queried during execution.

3. The kernel implementation should be such that the execution configuration (number of blocks and threads/block) affects the performance but not the results of the kernel invocation.

Methodology

NOTE: In this lab, you will not use SHARED MEMORY!

The program will be tested on the workstations and the CUDA server using the following matrix sizes:

1. Use C++ compiler (g++) to compile your code, and use new operator instead of malloc() to dynamically allocate memory (malloc may fail on very large memory allocations).

2. Try to design the data structures so to minimize the number of memory transactions between host and device (CPU and GPU).

3. DO NOT START TO CODE IMMEDIATELY. Spend some time designing your solution.

i. How do you handle matrix sizes exceeding the GPU capacity?

ii. How do you represent the matrices?

iii. Which memory transfers are involved with matrices within and beyond the GPU capacity?

iv. When designing the kernel, which work is performed by each thread? How do you correlate each thread with the data it processes?

4. Kernel calls use shared memory and registers. Your kernel should not use shared memory. To see how many registers are used by each thread, you can have a look at the GPU assembly file.The assembly file (called PTX file) can be generated by calling:

nvcc -ptx myfile.cu

This will generate myfile.ptx.

The PTX file will show you the assembly representation of your kernel. In particular, it will show you the code execute by each thread (as you know, all threads execute the same code!). The PTX file will include an area where the registers are declared. For example:

.reg .u16 %rh<4>;//16 bit registers

.reg .u32 %r<9>; //32 bit registers

.reg .u64 %rd<10>;//64 bit registers

.reg .pred %p<3>; // registers used for predication

If you know: how many registers are used by each thread and how many registers are available on the GPU, you can easily determine what is the maximum number of threads that you can run (for your particular kernel).

5. Use the occupancy calculator to calculate the optimal point for configuring the kernel.

Questions:

1. Run your kernel with different number of blocks and of threads/block and see how this affects the performance. Report GPU occupancy and execution time, and discuss the results. Consider execution configurations that are trivially bad, and compare them with good execution configurations. How do you know in advance that some configurations are "bad"? And what is a "good" execution configuration?

2. What is the largest execution configuration that you can use without exceeding the resources available on the GPU in use?