 
 
 	//Wind it down!
-	for (int k = 0; k>height; k++)
+	for (int k = 0; k<height; k++)
 	{
 		// This tells you if the current row is the first or second.
 		shft_wr = (k & 1);
 		// Read and write are opposite rows.
-		shft_rd = (THREADBLK+2)*((shft_wr+1) & 1);
+		shft_rd = base*((shft_wr+1) & 1);
 		//Block 0 is split so it needs a different algorithm.  This algorithm
 		//is slightly different than top triangle as described in the note above.
 
-		if (tid <= ((THREADBLK+1)-k) && tid >= (k-2))
+		if (tid < (THREADBLK-k) && tid > k)
 		{
 			temper[tidp + ((THREADBLK+2)*shft_wr)] = fo * (temper[tid+shft_rd] + temper[tid+shft_rd+2]) + (1.f-2.f*fo) * temper[tidp+shft_rd];
 		}
-
-	}
-
-
-
+	}
+
+	//Make sure the threads are synced
+	__syncthreads();
+
+	//Now thread 0 in each block (which never computes a value) is used to
+	//fill the shared right and left arrays with the relevant values.
+	//This grabs the top and bottom edges on the iteration when the top
+	//row is written.
+	if (shft_wr && tid < 4)
+	{
+		sL[k+itr+tid] = temper[(tid/2*(THREADBLK-1))+(tid-1)+k];
+		sR[k+itr+tid] = temper[((tid+2)/2*(THREADBLK-1))+(tid&1)-k];
+		itr += 2;
+	}
+
+	__syncthreads();
 
 }
 

__global__ void splitDiamond(REAL *right, REAL *left)
{

	__shared__ REAL temper[(2*THREADBLK)+4];
	__shared__ REAL sR[THREADBLK];
	__shared__ REAL sL[THREADBLK];



}


	// Pass to the left so all checks are for block 0 (this reduces arithmetic).
	// The left ridge is always kept by the block.
	sL[tid] = right[gid];

	// The right ridge is passed, each block 1-end gets the right of 0-end-1
	// Block 0 gets the right of the last block.
	if (blockIdx.x == (gridDim.x-1))
	{
		sR[tid] = left[tid];
	}
	else
	{
		sR[tid] = left[gid+blockDim.x];
	}

	// Initialize temper. Kind of an unrolled for loop.  This is actually at
