Un-coalesced global memory access caused by indirect access in CUDA

6
votes

My CUDA programm is suffering from un-coalesced global memory access. Although the idx-th thread only deal with the [idx]-th cell in an array, there are many indirect memory accesses as shown below. 

int idx=blockDim.x*blockIdx.x+threadIdx.x;

.... = FF[m_front[m_fside[idx]]];

For m_fisde[idx], we have coalesced accesses, but what we actually need is FF[m_front[m_fside[idx]]]. There is a two-level indirect access.

I tried to find some patterns of the data in m_front or m_fsied in order to make this to be a direct sequential access, but found out that they are almost 'random'.

Is there a possible way to tackle this?

1
cudagpugpgpu
This is effectively the same problem as sparse matrix addressing, and there has been rather a lot of work done on understanding how to improve that. You might get some ideas from looking at the literature on sparse matrix operations on GPUS. – talonmies
If there is any locality in the accesses, this question may be of interest. – Robert Crovella
@RobertCrovella... The Texture Mechanism link provided in the above linked answer, is expired. Can you please update the link? – sgarizvi
@sgar91 The link is fixed now. – Robert Crovella
@thierry Please, see my revised and improved answer. – Vitality

1 Answers

4
votes

Accelerating global memory random access: Invalidating the L1 cache line

Fermi and Kepler architectures support two types of loads from global memory. Full caching is the default mode, it attempts to hit in L1, then L2, then GMEM and the load granularity is 128-byte line. L2-only attempts to hit in L2, then GMEM and the load granularity is 32-bytes. For certain random access patterns, memory efficiency can be increased by invalidating L1 and exploiting the lower granularity of L2. This can be done by compiling with –Xptxas –dlcm=cg option to nvcc.

General guidelines for accelerating global memory access: disabling ECC support

Fermi and Kepler GPUs support Error Correcting Code (ECC), and ECC is enabled by default. ECC reduces peak memory bandwidth and is requested to enhance data integrity in applications like medical imaging and large-scale cluster computing. If not needed, it can be disabled for improved performance using the nvidia-smi utility on Linux (see the link), or via Control Panel on Microsoft Windows systems. Note that toggling ECC on or off requires a reboot to take effect.

General guidelines for accelerating global memory access on Kepler: using read-only data cache

Kepler features a 48KB cache for data that is known to be read‐only for the duration of the function. Use of the read‐only path is beneficial because it offloads the Shared/L1 cache path and it supports full speed unaligned memory access. Use of the read‐only path can be managed automatically by the compiler (use the const __restrict keyword) or explicitly (use the __ldg() intrinsic) by the programmer.