 	HANDLE_ERROR(cudaStatus);
 
 	// Allocate GPU buffers for matrix of shortest paths d(G)
-	cudaStatus =  cudaMalloc((void**)&dev_d, n * n * sizeof(int));
-	HANDLE_ERROR(cudaStatus);
-	cudaStatus =  cudaMalloc((void**)&dev_p, n * n * sizeof(int));
-	HANDLE_ERROR(cudaStatus);
+	cudaStatus = cudaMallocPitch(&dev_d, &pitch, n * sizeof(int), n);
+	HANDLE_ERROR(cudaStatus);
+
+	cudaStatus = cudaMallocPitch(&dev_p, &pitch, n * sizeof(int), n);
+	HANDLE_ERROR(cudaStatus);
+
+	pitch_int = pitch / sizeof(int);
 
         // Wake up gpu
 	wake_gpu_kernel<<<1, dimBlockP1>>>(32);
 	
         // Copy input vectors from host memory to GPU buffers.
-        cudaStatus = cudaMemcpyAsync(dev_d, G, n * n * sizeof(int), CMCPYHTD, cpyStream);
-	cudaStatus = cudaMemcpyAsync(dev_p, p, n * n * sizeof(int), CMCPYHTD, cpyStream);
-        
+        cudaMemcpy2DAsync(dev_d, pitch, G, n * sizeof(int), n * sizeof(int), n, CMCPYHTD, cpyStream);
+	cudaMemcpy2DAsync(dev_p, pitch, p, n * sizeof(int), n * sizeof(int), n, CMCPYHTD, cpyStream);
+
 	// cudaDeviceSynchronize waits for the kernel to finish, and returns
         cudaStatus = cudaDeviceSynchronize();
         HANDLE_ERROR(cudaStatus);
 	
 	FOR(block, 0, numOfBlock) 
 	{
-		fw_kernel_phase_one<VIRTUAL_BLOCK_SIZE><<<1, dimBlockP1>>>(block, n, dev_d, dev_p);
-		fw_kernel_phase_two<VIRTUAL_BLOCK_SIZE><<<dimGridP2, dimBlockP2>>>(block, n, dev_d, dev_p);
-		fw_kernel_phase_three<BLOCK_SIZE, THREAD_SIZE><<<dimGridP3, dimBlockP3>>>(block, n, dev_d, dev_p);
+		fw_kernel_phase_one<VIRTUAL_BLOCK_SIZE><<<1, dimBlockP1>>>(block, n, pitch_int, dev_d, dev_p);
+		fw_kernel_phase_two<VIRTUAL_BLOCK_SIZE><<<dimGridP2, dimBlockP2>>>(block, n, pitch_int, dev_d, dev_p);
+		fw_kernel_phase_three<BLOCK_SIZE, THREAD_SIZE><<<dimGridP3, dimBlockP3>>>(block, n, pitch_int, dev_d, dev_p);
 	}
 
 	// Check for any errors launching the kernel


	if (v1 < n && v2 < n)
        {
        	d[v1 * n + v2] = current_d[ty][tx];
		p[v1 * n + v2] = current_p[ty][tx];
        }
}


    	cudaStatus = cudaDeviceSynchronize();
	HANDLE_ERROR(cudaStatus);

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