adding back the BK check

and some template param tuning

A is laid out in column major in global mem

it's now outside the kernel, so we check only once instead of however
many times. However, this bloats binary size.

also changed template parameters a bit

since we know that As and Bs are laid out one after another in memory
(ie they are basically svals), and the overalls dimension is (BM + BN) *
BK, we just write one "nested" loop that does the zeroing

they're not used in the __global__ functions anyway

it's now a specialization of matmul_block2d