University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2023
Views:

491

Pages:

23
Author:

dugninisjfh1

Example - Optimizing a parallel reduction ‣ Slides from Mark Harris - NVIDIA ‣ Optimizing parallel reduction within a block - Tests over 4M entries - No global synchronization in CUDA - Recursive kernel invocation necessary to complete ‣ Reductions are a memory bound problem - Will measure optimization using bandwidth ‣ G80 GPU used in this example - 384 bit memory interface, 900MHz DDR - (384/8)*900*2/103 = 86.4GB/s Example - Optimizing a parallel reduction ‣ Interleaved addressing __global__ voidreduce0(int*g_idata, int*g_odata){ extern __shared__ int sdata[]; // each thread loadsone element from global to shared mem unsigned int tid = threadIdx.x; unsigned int i= blockIdx.x*blockDim.x+ threadIdx.x; sdata[tid] = g_idata[i]; __syncthreads(); // do reduction in shared mem for(unsigned ints=1; s < blockDim.x; s *= 2) { if(tid % (2*s) == 0){ sdata[tid] += sdata[tid + s]; } __syncthreads(); } // write result for this block to global mem if(tid == 0) g_odata[blockIdx.x] = sdata[0]; } Example - Optimizing a parallel reduction ‣ Interleaved addressing Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction ‣ Avoid divergent branching Change for(unsigned int s=1; s < blockDim.x; s *= 2) { if(tid % (2*s) == 0){ sdata[tid] += sdata[tid + s]; } __syncthreads(); } by for(unsigned int s=1; s < blockDim.x; s *= 2) { int index = 2 * s * tid; if (index < blockDim.x) { sdata[index] += sdata[index + s]; } __syncthreads(); } No divergence, but bank conflicts arise Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction ‣ Sequential addressing Example - Optimizing a parallel reduction ‣ Sequential addressing Change for(unsigned int s=1; s < blockDim.x; s *= 2) { int index = 2 * s * tid; if (index < blockDim.x) { sdata[index] += sdata[index + s]; } __syncthreads(); } by for(unsigned int s=blockDim.x/2; s>0; s>>=1) { if (tid < s) { sdata[tid] += sdata[tid + s]; } __syncthreads(); } Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction ‣ First add during load Change unsigned int tid = threadIdx.x; unsigned int i= blockIdx.x*blockDim.x+ threadIdx.x; sdata[tid] = g_idata[i]; __syncthreads(); by // perform first level of reduction, // reading from global memory, writing to shared memory unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; __syncthreads(); Example - Optimizing a parallel reduction ‣ Still far from peak bandwidth (86GB/s) - We can’t be bounded by bandwidth ‣ Look into instruction overhead (address arithmetic and loops) ‣ Let’s unroll loops Example - Optimizing a parallel reduction ‣ Unroll last warp - As instruction proceeds, number of active threads decreases - When we have less than 32 threads ‣ We don’t need __syncthreads() ‣ Don’t need if statement “if (tid < s)” Example - Optimizing a parallel reduction ‣ Unroll last warp for(unsigned int s=blockDim.x/2; s>32; s>>=1) { if (tid < s) sdata[tid] += sdata[tid + s]; __syncthreads(); } if (tid < 32) { sdata[tid] sdata[tid] sdata[tid] sdata[tid] sdata[tid] sdata[tid] } += += += += += += sdata[tid sdata[tid sdata[tid sdata[tid sdata[tid sdata[tid + 32]; + 16]; + 8]; + 4]; + 2]; + 1]; Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction ‣ Complete unrolling - If we knew the number of iterations at compile time we could completely unroll reduction ‣ Luckily for G80 we’re limited to 512 threads - Use templates! ‣ Some of the branching will be evaluated at compile time ‣ CUDA supports C++ templating on device and host functions - Block size will be our template parameter template __global__ void reduce5(int *g_idata, int *g_odata) Example - Optimizing a parallel reduction ‣ Complete unrolling if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; }__syncthreads(); } if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); } if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); } if (tid < 32){ if (blockSize >= 64) sdata[tid] += sdata[tid + 32]; if (blockSize >= 32) sdata[tid] += sdata[tid + 16]; if (blockSize >= 16) sdata[tid] += sdata[tid + 8]; if (blockSize >= 8) sdata[tid] += sdata[tid + 4]; if (blockSize >= 4) sdata[tid] += sdata[tid + 2]; if (blockSize >= 2) sdata[tid] += sdata[tid + 1]; Example - Optimizing a parallel reduction ‣ We don’t need to know the block size at compile time switch (threads) { case 512: reduce5<512><<< case 256: reduce5<256><<< case 128: reduce5<128><<< case 64: reduce5< 64><<< case 32: reduce5< 32><<< case 16: reduce5< 16><<< case 8: reduce5< 8><<< case 4: reduce5< 4><<< case 2: reduce5< 2><<< case 1: reduce5< 1><<< } dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; dimGrid, dimBlock, smemSize >>>(d_idata, d_odata); break; Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction ‣ Mix parallel and sequential execution to find optimal point - Instead of doing the first add when loading to shared memory, do as many as necessary Change unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; __syncthreads(); by unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x; unsigned int gridSize = blockSize*2*gridDim.x; sdata[tid] = 0; while (i < n){ sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; } __syncthreads(); Example - Optimizing a parallel reduction Example - Optimizing a parallel reduction template __global__ voidreduce6(int *g_idata, int *g_odata, unsigned int n) { extern __shared__ int sdata[]; unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*(blockSize*2) + tid; unsigned int gridSize = blockSize*2*gridDim.x; sdata[tid] = 0; do{sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; } while (i < n); __syncthreads(); if (blockSize >= 512) {if(tid<256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); } if (blockSize >= 256) {if(tid<128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); } if (blockSize >= 128) {if(tid< 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); } if (tid < 32){ if (blockSize >= 64) sdata[tid] += sdata[tid + 32]; if (blockSize >= 32) sdata[tid] += sdata[tid + 16]; if (blockSize >= 16) sdata[tid] += sdata[tid + 8]; if (blockSize >= 8) sdata[tid] += sdata[tid + 4]; if (blockSize >= 4) sdata[tid] += sdata[tid + 2]; if (blockSize >= 2) sdata[tid] += sdata[tid + 1]; } if (tid == 0) g_odata[blockIdx.x] = sdata[0]; } Example - Optimizing a parallel reduction Optimizations - Wrapping up ‣ Nice summary in CUDA Best Practices Guide Appendix A ‣ Basic strategies - Maximize parallel execution - Optimize memory usage to achieve maximum bandwidth - Optimize instruction usage to achieve maximum throughput

GPU Computing with CUDA - Optimizations Wrapping Up

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock