 * @param d matrix of shortest paths d(G)
 * @param p matrix of predecessors p(G)
 */
-template <int BLOCK_SIZE> __global__ void fw_kernel_phase_one(const unsigned int block, const unsigned int n, int * const d, int * const p)
+template <int BLOCK_SIZE> __global__ void fw_kernel_phase_one(const unsigned int block, const unsigned int n, const size_t pitch, int * const d, int * const p)
 {
 	int newPath;
 	int newPred;

    	cudaStatus = cudaDeviceSynchronize();
	HANDLE_ERROR(cudaStatus);

	cudaStatus = cudaMemcpy(d, dev_d, n * n * sizeof(int), CMCPYDTH);
	HANDLE_ERROR(cudaStatus);
	
	cudaStatus = cudaMemcpy(p, dev_p, n * n * sizeof(int), CMCPYDTH);
	HANDLE_ERROR(cudaStatus);
	
	cudaStatus = cudaFree(dev_d);

* @param p matrix of predecessors p(G)
*/

template <int BLOCK_SIZE, int THREAD_SIZE> __global__ void fw_kernel_phase_three(const unsigned int block, const unsigned int n, const size_t pitch, int * const d, int * const p)
{
	if (blockIdx.x == block || blockIdx.y == block) return ;
	int newPath;
