A single 640x640 matrix transposition should be large enough to have enough threads to keep a GPU busy. You don’t need new storage for all the matrices at once, you just need one extra set of 640x640 storage in order to not do an in-place transpose, and be able to use the efficient code described here:
[url]http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/[/url]
Take the first matrix, transpose it out-of-place to your new set of 640x640 storage. Then take the next matrix and transpose it out-of-place to the location occupied originally by the first matrix. Repeat this process for each matrix in sequence.
Even if you don’t like this approach, if you study the description given in the parallel forall blog, you will see that the matrix is transposed a tile at a time. For this square case, and since the tiles are also square, the source tile and destination tile effectively change places, so it should be possible to modify that algorithm to an in-place one, as long as 2 tiles (source and destination) are both loaded into shared memory, before any writes occur. Then two tile processing operations will be performed per threadblock (and only half as many threadblocks will be needed), and the operation can be done in-place.
Threadblocks on the main diagonal will require special-casing. Those tiles (blockIdx.x = blockIdx.y) have a source and destination tile location that is already the same, so they can be processed in a fashion identical to what is given in the parallel forall blog post. No need to load source and destination tiles before processing, they are the same.