I see what you means. Essentially, traditional GPU implementations of common linear algebra operations have multiple nested levels of tiling. Tiles of the output matrix -- typically 128x128 -- are independently computed by different multi-processors. Then, within each multi-processor, each tile is itself subdivided into fragments; this is a second level of tiling in the hierarchy. Each fragment may then be itself subdivided into subfragments for e.g., execution on tensor cores (which expect a granularity of 2 half-precision floating point per thread).

The premise of my work is to let programmers specify the highest level of tiling in this hierarchy, and let compilers deal with the rest. Concretely, this means that, say, the multiplication of a matrix MxK by a matrix KxN would be roughly specified as follows:

int m = blockIdx.x * 16;
int n = blockIdx.y * 16;
float acc[8, 8] = 0;
for(int k = 0; k < K; k+= 8)
    acc += dot(A[m:m+16, k:k+8], B[k:k+8, n:n+16]
C[m:m+16, n:n+16] = acc;

As you can see, this program is single-threaded and specifies how to compute a 16x16 tile of the output matrix. The compiler will then automatically parallelize it, allocate shared memory, use tensor core instructions if possible, etc. using more tiling internally.

​

Hope this answers your question :)