Ldmatrix access pattern to shared memory

Accelerated Computing CUDA CUDA Programming and Performance
1

hello

i’m trying to understand more the access patterns to shared memory when running a ldmatrix instruction (in order to mitigate bank conflicts)

the PTX docs says:
“When reading 8x8 matrices, a group of four consecutive threads loads 16 bytes.”, does it mean each of them reads 4 bytes? or a single thread does a vectorized load and then somehow communicate the data to the 3 others?.

Also the 8 groups (for the 8 rows) of four consecutive threads load their 16 corresponding bytes simultaneously ?

The figure 104 ( 1. Introduction — PTX ISA 9.3 documentation ) represents what each thread of the warp is going to hold after ldmatrix finishes, but is it actually the way the data is accessed as well ?

And if i’m using ldmatrix.x2, the second matrix load happens afterwards, they’re not happening in paralell right ?

sorry if those are obvious questions for you but i’m struggling to wrap my head around this

2

ldmatrix (without .x2) loads 128 bytes with one instruction, comprising 8 (aligned by 16 bytes) shared memory regions of 16 bytes.

To avoid bank conflicts (and therefore to have the reads simultaneously within one cycle) all those eight 16-byte regions (spanning 4 32-bit banks each) have to access distinctive shared memory banks (or use broadcast by accessing the identical 16-byte blocks).

It is not publicly or visibly defined, ‘what thread or lane does the read’. The warp does the reads as a unit. It is a .sync operation, so the lanes all participate.

The best mental model is that for ldmatrix providing the address by certain threads is not directly related to who receives the data from the address (it is defined in the PTX manual, but just because threads 0-7 provide addr0-7, not only threads 0-7 receive the results). And that the read is not done by certain threads, but overall by the ldmatrix instruction and that it also distributes the 16-bit values. See it as a unified instruction, not as two instructions (e.g. read+shuffle).

I have to admit, I am currently not sure (don’t remember) whether for ldmatrix.x2 addr0-addr7 have to be bank-conflict-free on their own, or whether all constituent addresses in addr0-addr7 and addr8-addr15 are resorted in an optimal way for accessing with the least amount of bank conflicts.

My private short note for ldmatrix.x4 (it would be similar for x1 and x2, just shortened to 2 or 4 addresses instead of 8):

// pointers have to be 16-Byte aligned
// 
// 0: 0, 8, 40, 48, ...
// 1: 10, 18, 50, 58, ...
// 2: 20, 28, 60, 68, ...
// 3: 30, 38, 70, 78, ...
// 4: 1, 9, 41, 49, ...
// ...
// 31: 37, 3F, 77, 7F, ...
//
// Thread 0 reads p[0][0], p[1][0], p[8][0], p[9][0], p[16][0], p[17][0], p[24][0], p[25][0]
// Thread 1 reads p[2][0], p[3][0], p[10][0], p[11][0], p[18][0], p[19][0], p[26][0], p[27][0]
// Thread 3 reads p[6][0], p[7][0], p[14][0], p[15][0], p[22][0], p[23][0], p[30][0], p[31][0]
// Thread 4 reads p[0][1], p[1][1], p[8][1], p[9][1], p[16][1], p[17][1], p[24][1], p[25][1]
// Thread 28 reads p[0][7], p[1][7], p[8][7], p[9][7], p[16][7], p[17][7], p[24][7], p[25][7]
// Thread 31 reads p[6][7], p[7][7], p[14][7], p[15][7], p[22][7], p[23][7], p[30][7], p[31][7]

The upper shown addresses (0, 8, 40, 48, …) are the flat hex format addresses, if the pointers are to consecutive 16-byte regions.
The lower shown addresses p[0][0], p[1][0] are if p[i] is the addr_i and [0..7] is the index into eight 16-bit values for the more flexible case of arbitrary (but aligned) pointers.

The best feature of ldmatrix is that it is possible to read 16-bit values per lane, i.e. split 32-bit values from shared memory. And still get the full bandwidth and no bank conflicts (if addresses and storage format are chosen in a clever way).