Global memory is served in 32-byte sectors, not individual words. Coalescing is arranging accesses so that the 32 threads of a warp touch as few sectors as possible, and use all of what they pull in.

out[i] = in[i];          // coalesced: warp reads 128 contiguous bytes, 4 sectors
out[i] = in[i * 16];     // strided: 32 separate sectors, 16x the traffic

Both lines move 128 useful bytes. The second one drags 1024 bytes across the bus to do it, and the wasted fraction never shows up as a stall in the profiler, only as a bandwidth figure well under spec.

The rule to aim for is that consecutive threadIdx.x values should touch consecutive addresses. This makes array-of-structs layouts a common problem, since the natural indexing gives every thread its own sector:

struct Particle { float x, y, z, vx, vy, vz; };
Particle *p;
p[i].x += 1.0f;          // stride of 24 bytes between neighbouring threads

Switching to a struct-of-arrays layout restores stride-1 access on each field:

struct Particles { float *x, *y, *z, *vx, *vy, *vz; };
p.x[i] += 1.0f;          // fully coalesced

Struct-of-arrays is the usual layout for GPU data for exactly this reason, even though it reads worse on the host side.

When the access pattern cannot be changed, for instance in a matrix transpose where one of the two accesses is strided by definition, the standard fix is to stage the tile through shared memory. Read coalesced, write to shared, synchronise, then read from shared in the awkward order and write coalesced. Shared memory has no coalescing requirement, only bank conflicts.