  * @param data: unique ptr to graph data with allocated fields on host
  */
 void cudaBlockedFW(const std::unique_ptr<graphAPSPTopology>& dataHost) {
-    // TODO not implemented yet
+    HANDLE_ERROR(cudaSetDevice(0));
+    int nvertex = dataHost->nvertex;
+    int *graphDevice, *predDevice;
+    size_t pitch = _cudaMoveMemoryToDevice(dataHost, &graphDevice, &predDevice);
+
+    dim3 gridDependedntPhase(1 ,1, 1);
+    dim3 blockDependentPhase(BLOCK_WIDTH, BLOCK_WIDTH, 1);
+
+    for(int round = 0; round < ; ++round) {
+        // Start dependent phase
+        _blocked_fw_dependent_ph<<<dimGrid, dimBlock>>>(block, pitch / sizeof(int), nvertex, graphDevice, predDevice);
+
+        // Start partially dependent phase
+
+        // Start independent phase
+    }
+
+    // Check for any errors launching the kernel
+    HANDLE_ERROR(cudaGetLastError());
+    HANDLE_ERROR(cudaDeviceSynchronize());
+    _cudaMoveMemoryToHost(graphDevice, predDevice, dataHost, pitch);
 }

            pred[indexYX] = pred[indexUX];
        }
    }
}

/**

    dim3 dimGrid((nvertex - 1) / BLOCK_WIDTH + 1, (nvertex - 1) / BLOCK_WIDTH + 1, 1);
    dim3 dimBlock(BLOCK_WIDTH, BLOCK_WIDTH, 1);

    int *graphDevice, *predDevice;
    size_t pitch = _cudaMoveMemoryToDevice(dataHost, &graphDevice, &predDevice);

    cudaFuncSetCacheConfig(_naive_fw_kernel, cudaFuncCachePreferL1);
    for(int vertex = 0; vertex < nvertex; ++vertex) {
        _naive_fw_kernel<<<dimGrid, dimBlock>>>(vertex, pitch / sizeof(int), nvertex, graphDevice, predDevice);
    }

    // Check for any errors launching the kernel
    HANDLE_ERROR(cudaGetLastError());
    HANDLE_ERROR(cudaDeviceSynchronize());
    _cudaMoveMemoryToHost(graphDevice, predDevice, dataHost, pitch);
}

