Memoryless Matrix Transposition II

- 4 mins read

Last time I presented a clever way of transposing that exploited the GPU’s support of 32-byte stores as its smallest “100% efficient” transaction size.

An open question was whether designing a transpose that performed 64-byte stores would achieve better performance at the cost of more instructions.

The answer is yes. Performing 64-byte transactions improves throughput by ~8%. The new 64-byte transpose kernel reaches ~148 GB/sec. on a 1024x1024 matrix of 32-bit elements on a K20c (758MHz). This is getting very close to the ~153 GB/sec. demonstrated by the shared transpose example.

Tesla K20c : sm_35 * 13 transposeKernel<<<1024,64>>>(1024 x 1024)
.... Validated! loops (100) : avg 0.05274 ms. = 148.12 GB/sec.

The interesting difference between this approach and last time is that of the different approaches I tried, the highest performing kernel dispensed with SHFL’s and implicitly performed rotations by manipulating the lower bits of the load and store pointers.

Similar to last time, neither __syncthreads() calls or shared memory are required.

Here is an illustration of how I chose to rearrange a tile of 4 lanes each holding 16 32-bit values:

A quick explanation:

Each lane loads 16 values. A logical tile contains 16x4 elements and occupies 4 lanes. I explicitly show the rotations that are performed on load but do not show the counter-rotations that must be performed on store. More on that later.

The rotations can be performed by either SHFL ops or implicitly through pointer manipulation ("swizzling on load"). Implicit “rotation pointers” are created by add-masking or using a BFI instruction. On load, each 4x4 block of registers is rotated one lane to the right of the block directly above it. With either approach, the warp performs 16 standard coalesced load transactions.

The exchange phase simply rearranges values in lanes 2 and 3 of the tile so that the next phase can use the SLCT opcode to choose which value to store back to device memory. More simply, this is every warp lane id that has bit 2 enabled.

Finally, the select phase takes advantage of the symmetries created by the rotation and exchange phases to enable all 4 lanes to coordinate a 64-byte (16-element) write. If you inspect each set of same-colored blocks you will see that they can be selected and written as vec4 instances simply by checking whether the lane is odd or even.

The 16x4 registers are written by the 4 coordinating lanes as 4 vec4 quads in the order: red, orange, green, blue. They’re each separated by a matrix width of elements.

As noted above, the 4x4 blocks of registers across the 4 lanes can either be counter-rotated using SHFL or the store pointer can be manipulated to perform the same operation implicitly ("swizzling on store").

That’s it and the performance is pretty good.

I found that the pointer manipulation approach was slightly faster than using SHFL ops. Conversely, in the previous blog post the SHFL implementation was slightly faster than the pointer approach when performing 32-byte stores.

There are quite a few ways to avoid using shared memory but… it’s probably not possible to get that last 10 GB/sec. of throughput achieved by NVIDIA’s shared transpose example on 1024x1024x32-bit matrices. The number of instructions in that kernel is tiny.

But if you’re interested in creating your own specialized transpose or unique data-rearrangement kernel, here is a list of low-level GPU features that you can mix-and-match:

  • writing 16-byte v2 or v4 types to simplify inter-lane word movement
  • explicit shuffling
  • rotations via shuffling
  • implicit shuffling of vector components (swizzling) on load
  • implicit shuffling of vector components (swizzling) on store
  • maximizing use of the SELP opcode: d = pred ? a : b;
  • thinking in terms of the non-existent opcode MOVP: if (pred) a = c; else b = c;
  • exchange two registers with an XCHG: t = a; a = b; b = t;

Please let me know if you have any questions or feedback.