         ix = thread_index % ifaces_x;
         iy = thread_index/ifaces_y;
 
-        I = (ix + mbc)*xs + (iy + mbc)*ys;
-
-        if (ix < mx && iy < my)   
+        I = (iy + mbc)*ys + (ix + mbc);
+
+        //if (ix < mx && iy < my)   
         {
 
             for(mq = 0; mq < meqn; mq++)

        ix = thread_index % ifaces_x;
        iy = thread_index/ifaces_y;

        I = (ix + mbc)*xs + (iy + mbc)*ys;

        if (ix < mx && 0 < iy && iy < my+1)   /* Is this needed? */
        {

            for(mq = 0; mq < meqn; mq++)


    maxcfl = 0;

    ifaces_x = mx + 2*mbc-1;
    ifaces_y = my + 2*mbc-1;
    num_ifaces = ifaces_x*ifaces_y;

    if (b4step2 != NULL)
    {
        ifaces_x = mx + 2*mbc - 1;           /* Visit all interior left edge */
        ifaces_y = my + 2*mbc - 1;           /* Visit all interior bottom edges */
        num_ifaces = ifaces_x*ifaces_y;

        for(thread_index = threadIdx.x; thread_index < num_ifaces; thread_index += blockDim.x)
        {
            ix = thread_index % ifaces_x;
            iy = thread_index/ifaces_y;

            I = (iy + 1)*ys + (ix + 1)*xs;  /* Start at one cell from left/bottom */

            //if (ix < mx + 2*mbc-1 && iy < my + 2*mbc-1)
            {
                for(mq = 0; mq < meqn; mq++)
                {
