a Goal:
Design program that computes square matrix multiplication on GPU using CUDA. 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.
NOTE: In this lab, you will not use SHARED MEMORY!
The program will be tested on the workstations and the cuda1 server using the following matrix
sizes:
Hints:
1. Have a look at CUDA library reference website at:
http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/online/
You will use the CUDA Runtime API:
http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/online/group__C
UDART.html
In particular, the Device Management module contains functions that allow you to get
information about the hardware resources of the GPU in the system.
2. Requirement #3 implies that the sizes of the matrices must be passed as parameters to
your matrix multiplication kernel. In other word, Cm,k = Am,n * Bn,k can be implemented
by a kernel with the following signature:
void matmul(float *a, float *b, float *c, unsigned m, unsigned n, unsigned k);
3. 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).
4. Try to design the data structures so to minimize the number of memory transactions
between host and device (CPU and GPU).
5. 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 within matrices 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?
6. 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).
7. Use the occupancy calculator to calculate the optimal point for configuring the kernel.
The occupancy calculator can be downloaded from:
http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls
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?