 		cudaErr(cudaMemcpy(sobelYFilter, hostSobelYFilter, 9 * sizeof(float), cudaMemcpyHostToDevice));
 	}
 
-	static GPUFrame sobelX(image.width(), image.height(), 4 * image.width(), image.height(), image.timestamp());
-	static GPUFrame sobelY(image.width(), image.height(), 4 * image.width(), image.height(), image.timestamp());
+	static GPUFrame sobelX(image.width(), image.height(), 4 * image.width(), image.height(), 0);
+	static GPUFrame sobelY(image.width(), image.height(), 4 * image.width(), image.height(), 0);
 
 	// figure out dimensions
-	dim3 grid, block(BLOCK_HEIGHT, BLOCK_WIDTH);
-	grid.x = image.width() / (2 * block.x);
+	dim3 grid, block(BLOCK_WIDTH, BLOCK_HEIGHT);
+	grid.x = image.width() / (1 * block.x);
 	grid.y = image.height() / block.y;
 
-	if(image.width() % (2 * block.x))
+	if(image.width() % (1 * block.x))
 		grid.x++;
 
 	if(image.height() % block.y)
 		grid.y++;
 
-	std::cout << "grid: " << grid.x * block.x << " x " << grid.y * block.y << " threads" << std::endl;
-
 	// launch convolution kernel with sobel matrix
-	kernelMatrixConvolution<<< grid, block >>>(image.data(), image.pitch(), sobelX.data(), sobelX.pitch(), image.width(), image.height(), sobelXFilter, 3, 3);
-	kernelMatrixConvolution<<< grid, block >>>(image.data(), image.pitch(), sobelY.data(), sobelY.pitch(), image.width(), image.height(), sobelYFilter, 3, 3);
+	kernelMatrixConvolution<<< grid, block >>>(image.data(), image.pitch(),
+											   sobelX.data(), sobelX.pitch(),
+											   image.width(), image.height(),
+											   sobelXFilter,
+											   3, 3);
+
+	kernelMatrixConvolution<<< grid, block >>>(image.data(), image.pitch(),
+											   sobelY.data(), sobelY.pitch(),
+											   image.width(), image.height(),
+											   sobelYFilter,
+											   3, 3);
 
 	// vector sum of both sobel images
-	kernelVectorSum<<< grid, block >>>(sobelX.data(), sobelX.pitch(), sobelY.data(), sobelY.pitch(), edges.data(), edges.pitch(), image.width(), image.height());
+	kernelVectorSum<<< grid, block >>>(sobelX.data(), sobelX.pitch(),
+									   sobelY.data(), sobelY.pitch(),
+									   edges.data(), edges.pitch(),
+									   image.width(), image.height());
 
 	// sync and check for errors
 	cudaDeviceSynchronize(); cudaErr(cudaGetLastError());

}

// sum of magnitudes of A and B treated as legs of a right triangle
__global__
void kernelVectorSum(const void* const inputA, const unsigned pitchInputA,
					 const void* const inputB, const unsigned pitchInputB,

	const unsigned gridYidx = blockIdx.y * blockDim.y + threadIdx.y;

	// kill threads that are out of bounds
	if(/*gridXidx * 2 >= pixelsWidth || */gridYidx >= pixelsHeight)
		return;

	// destinations for packed data
