Abstract This blog will try to improve the performance of previous reduction algorithm by freeing strides.
1. Strides The existence of strides directly results in a discount of load or store efficiency. This is because half the threads in the transaction are not used and represent wasted bandwidth. Therefore, ensuring that as much as possible of the data fetched without strides is an important part of performance optimization of memory accesses.
2. Key Codes
// reduction for(int i = 1024/2; i > 0; i >>=1){ if(tid < i){ data[tid] += data[tid + i]; } __syncthreads(); }A operation diagram is presented to explain these codes figuratively as follows.
3. Experimental Results
The experimental Results shows a much higher performance achieved by freeing strides, which is a fundamental but indispensable strategy when optimizing CUDA codes.
4. More The source code can be viewed on Github.
