[system-track][computing][cuda]cuda : performance issue

memory coalescing. about global memory

https://devblogs.nvidia.com/how-access-global-memory-efficiently-cuda-c-kernels/

nvidia-gpu memory hierarchy:
…

coalesce global memory load/store of threads in a warp into transactions as small as possible, better utilize memory bandwidth

32-byte, 64-byte, 128-byte / transaction

• dis-aligned memory access e.g. a[thread.x+32] -> coalesce access into L1 cache 128-byte cache lines-> good bandwidth utilization
• stride memory access e.g. a[thread.x*32] -> bad locality, hard to coalesce -> low bandwidth utilization

shared memory. about on-chip memory

https://devblogs.nvidia.com/using-shared-memory-cuda-cc/
stride memory access, bad memory bandwidth utilization, e.g. access multi-dim array -> tile, shared mem

shared memory, much lower latency than global memory. but bank conflicts happen
shared memory allocated per thread block, shared within thread block
e.g.
__shared__ int s[64]; //static, , use shared memory explicitly
extern __shared__ int s[]; //dynamically

usage(combined with thread synchronization):

• user-managed data caches
• high-performance cooperative parallel algorithms, e.g. parallel reduction
• facilitate global memory coalescing in case not be possible without shared memory

threads need synchronization for shared memory access

• __syncthreads(), a barrier across a thread block

performance issue: bank conflict
how memory addresses map to memory banks:

• 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 2.0, the warp size is 32 threads and the number of banks is also 32. A shared memory request for a warp is not split as with devices of compute capability 1.x, meaning that bank conflicts can occur between threads in the first half of a warp and threads in the second half of the same warp.
• Devices of compute capability 3.x have configurable bank size, which can be set using cudaDeviceSetSharedMemConfig() to either four bytes (cudaSharedMemBankSizeFourByte, the default) or eight bytes (cudaSharedMemBankSizeEightByte). Setting the bank size to eight bytes can help avoid shared memory bank conflicts when accessing double precision data.

e.g. for usage 3, matrix transpose optimizing
https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/
/*
This mapping is up to the programmer; the important thing to remember is that to ensure memory coalescing we want to map the quickest varying component to contiguous elements in memory. In Fortran contiguous addresses correspond to the first index of a multidimensional array, and threadIdx.x and blockIdx.x vary quickest within blocks and grids, respectively.
*/

branch divergence / threads divergence

https://stackoverflow.com/questions/17223640/is-branch-divergence-really-so-bad
e.g. reduction

• naive code

__global__ void reduce(int *input, unsigned int N, int *total){
	unsigned int tid = threadIdx.x;
	unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
	__shared__ int x[BSIZE];
	x[tid] = (i<N) ? input[i] : 0;
	__syncthreads();

	for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
		__syncthreads();
		if (tid % (2*stride) == 0){
			x[tid] += x[tid + stride];
		}
	}
	if (tid == 0) atomicAdd(total,x[tid]);
}

• improved code : reduce divergence and bank conflict
  no divergence until stride < 32 ; all warps active when stride >=32

for (stride = blockDim.x/2; stride > 1; stride /= 2) { 
	 __syncthreads();
	if (tid < stride){
		x[tid] += x[tid + stride];
	}
}

tutorial from

https://devblogs.nvidia.com/even-easier-introduction-cuda/

https://www.google.com.hk/url?sa=t&rct=j&q=&esrc=s&source=web&cd=1&cad=rja&uact=8&ved=2ahUKEwj_ja-uy_zgAhWZOnAKHYfnAdYQFjAAegQIAxAC&url=http%3A%2F%2Fcseweb.ucsd.edu%2Fclasses%2Ffa12%2Fcse260-b%2FLectures%2FLec09.pdf&usg=AOvVaw2QqIAMqGhA6sfoxz_DN4Du

https://www.google.com.hk/url?sa=t&rct=j&q=&esrc=s&source=web&cd=2&cad=rja&uact=8&ved=2ahUKEwjy_7i6y_zgAhUDF4gKHXuADqQQFjABegQIBBAC&url=http%3A%2F%2Fhomepages.math.uic.edu%2F~jan%2Fmcs572%2Fmemory_coalescing.pdf&usg=AOvVaw2GXJmYi5FaJuOc9om5vC4w