Coalesced global memory

What a warp asks for

32 consecutive threads reading consecutive 4-byte words touch 128 contiguous bytes: four 32-byte sectors and one cache line. The memory system issues one request per line, so the warp receives 128 useful bytes for 128 bytes fetched. That is the whole of coalescing, and it is a property of one instruction across one warp.

Strided access buys sectors you do not use

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
c[row * N + col] = a[col * N + row];   // lanes walk a column of a

Here threadIdx.x walks columns of a with stride N. Every lane lands in a different sector, so the warp fetches 32 sectors to deliver 128 bytes. Unrolling the loop does not help; the addresses still miss each other.

Fix the layout or stage it

Either store the data in the order the kernel reads it, or load a tile coalesced into shared memory and read the transpose from there, where a stride costs at most a bank conflict and no DRAM traffic. Then measure: ncu reports sectors per request, and a fully coalesced 32-bit load sits at four. Metric names move between releases, so confirm l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum against the Profiling Guide for your version.

Things that break

A pointer offset by one float misaligns every access, so a 128-byte request spans two lines and the warp pays an extra sector per request; allocate aligned and index from zero. An array of 12-byte structs turns a unit-stride walk into a 12-byte stride. And a coalesced read followed by a strided write is still half a problem, which is why the transpose needs the tile too.