       }
     }
     else {
-//       typedef char x[(bx * by % mb == 0) ? 1 : -1];     // bx * by must be a multiple of mb
-//       typedef char y[((bx * by * 16) / mb == nb) ? 1 : -1];     // when the threads are wrapped around mb they must spread along to nb
+      // Read A from shared memory
 #pragma unroll
       for (int l = 0; l < kb; l++)
-        saxpy(a[(bx * by == mb) ? ti : ti % mb][l],
-              &b[l][(bx * by == mb) ? 0 : 16 * (ti / mb)], c);
+        saxpy(a[ti % mb][l], &b[l][16 * (ti / mb)], c);
     }
 
     __syncthreads();

 * kb is chosen to be the largest multiple of 16 such that the number of blocks
 * per multiprocessor is limited by the register usage.
 */
template void ssyrk<CBlasUpper, CBlasNoTrans, 64, 16, 16, 16,  4>(int, int, float, const float *, int, float, float *, int);
template void ssyrk<CBlasLower, CBlasNoTrans, 64, 16, 16, 16,  4>(int, int, float, const float *, int, float, float *, int);
template void ssyrk<CBlasUpper, CBlasTrans,   32, 32,  8,  8,  8>(int, int, float, const float *, int, float, float *, int);
template void ssyrk<CBlasLower, CBlasTrans,   32, 32,  8,  8,  8>(int, int, float, const float *, int, float, float *, int);


  while (k > 0) {
    if (trans == CBlasNoTrans) {
      // C = aAA' + bC so read B into shared memory and transpose leaving A
      // untransposed in global memory
#pragma unroll
      for (int l = 0; l < kb; l += by) {
#pragma unroll
