---
slideOptions:
transition: slide
---
# GPU
## Cuda intro
### GPU introduction
* memory: global --> PBSM (per block shared memory) --> local register
* multiple stream multi-processors (SM):
* vector machine, share register file, programmable cache
### Cuda
* Kernel = Many Concurrent Threads
* kernel = grid of thread blocks
* Each thread executes the same code but **on the different data based on its threadID**
* threads are grouped into thread blocks
* blocks are independent to each other
* Thread Level Scheduling - Warp
* 
* Once a thread block is scheduled to an SM, threads in the thread block are further partitioned into warps.
* warp consists of 32 consecutive threads and all threads in a warp are executed in SIMT ie all threads exe the same instructions but on its own private data
* Threads in a warp will be executing the same instruction
* Instructions are issued per warp
> Keep the number of threads per block a multiple of warp size (32).
> Avoid small block sizes: Start with at least 128 or 256 threads per block.
> Adjust block size up or down according to kernel resource requirements.
### cuda language
* `kernelFunc<<< nB, nT, nS, Sid >>>(...);`
* `nB`: number of block per grid
* `nT`: number of threads per block
* `nS`: shared memory size
* `threadIdx; blockIdx; blockDim; gridDim`
* eg1
```cuda
dim3 grid(3,2):
dim3 blk(5,3);
my_kernel<<< grid, blk>>>();
// (x,y) or (x,y,z)
```
* eg2
```cuda
// Kernel definition
__global__ void VecAdd(float* A)
{
int i = threadIdx.x * blockDim.x + threadIdx.y;
A[i] = A[i] + 1;
}
//...
size=10; dim3 blk(size, size);
my_kernel<<< 1, blk >>>(A, size);
```
* blockDim.x: number of threads per block
*
* function qualifier:
* `__device__`: Executed on the device, Callable from the device only
* `__global__ `: Executed on the device, **Callable from the host only (must have void return type!)**
* Variable qualifier:
* `__device__`: device’s global memory (shared among SMs) space
* `__constant__`
* `__shared__`: in **thread block's** shared memory
* Program compilation
* 
* Kernel Time Measurement
```cuda
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
kernel<<<block, thread>>>();
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float time;
cudaEventElapsedTime(&time, start, stop);
```
* Dynamic Parallelism
* launch new grids from the GPU
* http://on-demand.gputechconf.com/gtc/2012/presentations/S0338-New-Features-in-the-CUDA-Programming-Model.pdf
## GPU Architecture
### Thread execution
* execution model
* Hardware schedules thread blocks onto available SMs
* warp
* At any time, only one warp is executed per SM
* All threads in a Warp execute the same instruction when selected
* Threads in a wrap execute physically in parallel
* Warps and blocks execute logically in paralle
### memory hierarchy
* registers
* consumes zero extra clock cycles per instruction, except
* Register memory bank conflicts
* not indexable
* **Array variables** always are allocated in **local memory**
* Register spilling: Local memory is used if the register limit is met
* limit: cuda6.1:
* 65K registers per block
* 255 registers per thread
* local memory
* Usually when one runs out of SM resources
* "local" for each thread has its own private area
* **Array variables** always are allocated in **local memory**
* Stores are cached in L1
* Not really a “memory” – bytes are stored in global memory (DRAM)
* eg.
```
__device__ void distance(int m, int n, int *V){
int i, j, k; // register
int a[10], b[10], c[10]; // local memory, off-chip DRAM
...
}
```
* shared memory
* shared by all the **threads in a block**
* Size: at most 48K per block (CUDA 6.1)
* `extern __shared__ <type> <name>[]` (pg.7)
* dyanmic allocation: `nS` in `func<<<nB, nT, nS>>>(..);`
```cuda
extern __shared__ int S[][];
__global__ void FW_APSP(int k, int D[n][n]) {
int i = threadIdx.x;
int j = threadIdx.y;
S[i][j]=D[i][j]; // move data to shared memory
__syncthreads();
// do computation
if (S[i][j]>S[i][k]+S[k][j])
D[i][j]= S[i][k]+S[k][j];
}
```
* architecture
* Each bank can service one address per cycle
* Shared memory is as fast as register if no bank conflict
* Shared memory banks are organized such that successive 32-bit words are assigned to successive banks and the bandwidth is 32 bits per bank per clock cycle.
* For devices of compute capability 1.x, the warp size is 32 threads and the number of banks is 16
* https://stackoverflow.com/questions/7903566/how-is-2d-shared-memory-arranged-in-cuda
* global memory
* memory coalescing
* Global memory access happens in transactions of 32 or 128 bytes
* The hardware will try to reduce to as few transactions as possible
* Coalesced access:
> A group of 32 contiguous threads ("warp") accessing adjacent words. Few transactions and high utilization
* constant memory
* Same usage and scope as the global memory except
* Read only
* Declare by variable qualifier `__constant__`
* Move by cudaMemcpyToSymbol() & cudaMemcpyFromSymbol()
* warp divergence
## share memory
* ref: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-5-x
* For compute capability 5.x and 6.x:
* Shared memory has 32 banks that are organized such that successive **32-bit words** map to successive banks. Each bank has a bandwidth of 32 bits per clock cycle.
* bank index = (byte address / 4 bytes per bank) % 32 banks
## Reference
* https://docs.nvidia.com/cuda/
* https://nanxiao.me/en/about/
* http://www.icl.utk.edu/~luszczek/teaching/courses/fall2016/cosc462/pdf/GPU_Fundamentals.pdf
* https://medium.com/@smallfishbigsea/basic-concepts-in-gpu-computing-3388710e9239
* https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf
* http://on-demand.gputechconf.com/gtc-express/2011/presentations/cuda_webinars_WarpsAndOccupancy.pdf
* https://devblogs.nvidia.com/cuda-pro-tip-nvprof-your-handy-universal-gpu-profiler/
* http://www.enseignement.polytechnique.fr/profs/informatique/Eric.Goubault/Cours09/CUDA/SC07_CUDA_5_Optimization_Harris.pdf
* http://on-demand.gputechconf.com/gtc-express/2011/presentations/StreamsAndConcurrencyWebinar.pdf
* http://on-demand.gputechconf.com/gtc-express/2011/presentations/NVIDIA_GPU_Computing_Webinars_CUDA_Memory_Optimization.pdf
* http://on-demand.gputechconf.com/gtc-express/2011/presentations/GTC_Express_Sarah_Tariq_June2011.pdf
*
Sign in with Wallet
Connect another wallet