WarpTiling

@siboehm Thank you, this is still a great resource which I've been meaning to work through for ages and finally had time to. I've always wondered why warp tiling (forcing a specific tile to be processed by a single warp) would be better than processing the whole tile with all the threads. I've read your [summary of advantages](https://siboehm.com/articles/22/CUDA-MMM) but I still can't understand with the exception of a special "register cache on recent GPUs" how any of the fundamental features of how warps operate would have any impact when using one approach over the other. I can understand the advantage if the algorithm can make use of warp communication primatives or if each warp could work on an independant piece of global memory (avoid `__syncthreads()`) but I can't see why it would be advantageous in general so I think I'm missing something? Anyway taking your implementation and hyperparameters if I switch from warp tiling where each of the `4`, `64x64` warp tiles is processed by a single warp to [block tiling](https://github.com/cudawarped/SGEMM_CUDA/blob/blockSubTiling/src/kernels/6a_kernel_blocktiling.cuh) (or block sub tiling) where each of the `4`, `64x64` sub blocks is processed by all 128 threads then I get the roughly the same result for both approaches on an A6000, which is what I would have expected all else being equal (which it never is once nvcc is envolved). Results for block tilining apprach on an A6000 can be found [here](https://github.com/cudawarped/SGEMM_CUDA/tree/blockSubTiling). I know its been a few years since you wrote this but I was wondering if you had any thoughts on the observations above? Do you think warp tiling is a better approach on more modern GPU's?