in src/gpu_hmma.h [313:327]
__device__ __forceinline__ void mma_sync(fragmentA<OP_A,TILE> &fragA, fragmentB<OP_B,TILE> &fragB)
{
if (TILE == M16N16K16 || TILE == M8N8K16)
{
if (OP_A == OP_N && OP_B == OP_N) MMA_SYNC("m16n16k16", "row", "row");
if (OP_A == OP_T && OP_B == OP_N) MMA_SYNC("m16n16k16", "col", "row");
if (OP_A == OP_N && OP_B == OP_T) MMA_SYNC("m16n16k16", "row", "col");
}
else if (TILE == M8N32K16)
{
// m8n32k16 flag doesn't change mma behavior and isn't backwards compatable with cuda 9.0
if (OP_A == OP_N && OP_B == OP_N) MMA_SYNC("m16n16k16", "row", "row"); // m8n32k16
if (OP_A == OP_T && OP_B == OP_N) MMA_SYNC("m16n16k16", "col", "row"); // m8n32k16
}
}