 	__shared__ REAL sR[32];
 	__shared__ REAL sL[32];
 
-	int gid = blockDim.x * blockIdx.x + threadIdx.x; //Global Thread ID
+	//Same as upTriangle
+	int gid = blockDim.x * blockIdx.x + threadIdx.x;
 	int tid = threadIdx.x;
 	int shft_rd;
 	int shft_wr;
 
-	// Pass to the left so all checks are for block 0.
-	// The left ridge is kept by the block.
+	// Pass to the left so all checks are for block 0 (this reduces arithmetic).
+	// The left ridge is always kept by the block.
 	sR[tid] = left[gid];
 
 	// The right ridge is passed, each block 1-end gets the right of 0-end-1


	}

	right[gid] = sR[tid];
	left[gid] = sL[tid];

}


__global__ void downTriangle(REAL *IC, REAL *right, REAL *left)
{

#define REAL       float
#define TH_DIFF    8.418e-5

// Declare constant Fourier number that will go in Device constant memory.
__constant__ REAL fo;

/*
First function, the right side up triangle.  It takes IC which is the array
of temperatures at a single time and returns left and right values that are
either passed to the neighborning node or kept by the home node.  Elements of
both functions were inspired by the nvidia presentation on reduction here:
https://docs.nvidia.com/cuda/samples/6_Advanced/reduction/doc/reduction.pdf
*/

__global__ void upTriangle(REAL *IC, REAL *right, REAL *left)
{
	/*
	Initialize shared variables.  Each node (warp) will store 32 values on the
	right and left sides of their triangle, 2 on each side for each timeslice.
	Since the base of the triangle is 32 numbers for each node, 16 timeslices
	are evaluated per kernel call.
	Temper stores the temperatures at each timeslice.  Since only the current
	and previous timeslice results need to be held at each iteration.  This
	variable has 64 values, or two rows of 32, linearized.  The current and
	previous value alternate rows at each timeslice.
	*/
	__shared__ REAL temper[64];
	__shared__ REAL sR[32];
	__shared__ REAL sL[32];

	int gid = blockDim.x * blockIdx.x + threadIdx.x; //Global Thread ID
	int tid = threadIdx.x; //Warp or node ID
	int shft_wr; //Initialize the shift to the written row of temper.
	int shft_rd; //Initialize the shift to the read row (opposite of written)

	//Assign the initial values to the first row in temper, each warp (in this
	//case each block) has it's own version of temper shared among its threads.
	temper[tid] = IC[gid];
	__syncthreads(); // Then make sure each block of threads are synced.

	//This counter facilitates the transfer of the relevant values into the
	//right and left saved arrays.
	int itr = -1;

	//The initial conditions are timslice 0 so start k at 1.
	for (int k = 1; k<16; k++)
	{
		//Bitwise even odd. On even iterations write to first row.
		shft_wr = (k & 1);
		//On even iterations write to second row (starts at element 32)
		shft_rd = 32*((shft_wr+1) & 1);

		//Each iteration the triangle narrows.  When k = 1, 30 points are
		//computed, k = 2, 28 points.
		if (tid <= (31-k) && tid >= k)
		{
			temper[tid + (32*shft_wr)] = fo * (temper[tid+shft_rd-1] + temper[tid+shft_rd+1]) + (1-2.*fo) * temper[tid+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 == 0)
		{
			sL[k+itr] = temper[k-1];
			sL[k+itr+1] = temper[k];
