diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 048e6814..029725a5 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -5,12 +5,20 @@ on: jobs: build: - runs-on: ubuntu-latest + runs-on: ubuntu-20.04 steps: - uses: actions/checkout@v3 with: submodules: recursive + - name: Install CUDA + run: | + wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/cuda-keyring_1.0-1_all.deb + sudo dpkg -i cuda-keyring_1.0-1_all.deb + sudo apt-get update + sudo apt-get install -y cuda + echo "/usr/local/cuda-12.1/bin" >> $GITHUB_PATH + - name: Install prerequisites with apt run : | sudo apt-get update @@ -42,7 +50,7 @@ jobs: cd app mkdir build cd build - cmake .. -DCMAKE_BUILD_TYPE=Release + cmake .. -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_ARCHITECTURES=86 cmake --build . - name: Test diff --git a/Dockerfile b/Dockerfile index be2ca9af..a5eb77f7 100644 --- a/Dockerfile +++ b/Dockerfile @@ -1,5 +1,6 @@ FROM ubuntu:20.04 AS dev +ENV DEBIAN_FRONTEND=noninteractive ENV TZ=Europe/Lisbon RUN ln -snf /usr/share/zoneinfo/$TZ /etc/localtime && echo $TZ > /etc/timezone RUN apt-get update @@ -8,7 +9,16 @@ RUN apt-get install -y \ cmake \ git \ npm \ - python3 + python3 \ + wget + +## Add CUDA +WORKDIR /tmp/ +RUN wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/cuda-keyring_1.0-1_all.deb +RUN dpkg -i cuda-keyring_1.0-1_all.deb +RUN apt-get update +RUN apt-get install -y cuda +ENV PATH=/usr/local/cuda-12.1/bin${PATH:+:${PATH}} ## Configure Apache2 RUN a2enmod rewrite diff --git a/app/CMakeLists.txt b/app/CMakeLists.txt index 83a69c3a..84fbf671 100644 --- a/app/CMakeLists.txt +++ b/app/CMakeLists.txt @@ -1,6 +1,6 @@ cmake_minimum_required(VERSION 3.13) -project(dynaminator) +project(dynaminator LANGUAGES CXX CUDA) set(CMAKE_CXX_STANDARD 17) SET(CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS} "-std=c++17 \ @@ -27,8 +27,12 @@ SET(CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS} "-std=c++17 \ include_directories("include" "test/include" "script/include" "libs/rapidxml" "libs/CTPL" "libs/http-status-codes-cpp") # ################ SIMULATOR ################ -file(GLOB_RECURSE SRC "src/*.cpp" "*/src/*.cpp") +file(GLOB_RECURSE SRC "src/*.cpp" "*/src/*.cpp" "src/*.cu" "*/src/*.cu") add_executable(dynaminator main.cpp ${SRC}) +target_compile_options(dynaminator PRIVATE $<$: + -g + -G + >) # ################ TESTS ################ file(GLOB_RECURSE TESTS_SRC "test/main.cpp" "test/test_*.cpp" "test/src/*.cpp" "test/src/*/*.cpp") @@ -38,6 +42,10 @@ target_link_libraries(tests PRIVATE Catch2::Catch2WithMain stdc++fs) include(CTest) include(Catch) catch_discover_tests(tests) +target_compile_options(tests PRIVATE $<$: + -g + -G + >) # ################ SCRIPT ################ find_package(nlohmann_json 3.2.0 REQUIRED) diff --git a/app/include/Graph.hpp b/app/include/Graph.hpp index b821b9be..cc896d9b 100644 --- a/app/include/Graph.hpp +++ b/app/include/Graph.hpp @@ -2,17 +2,19 @@ #include #include +#include +#include class Graph { public: - typedef long Node; + typedef int32_t Node; static const Node NODE_INVALID = -1; struct Edge { - typedef double Weight; - typedef long ID; + typedef float Weight; + typedef int32_t ID; static const Weight WEIGHT_INF; diff --git a/app/include/cuda.hpp b/app/include/cuda.hpp new file mode 100644 index 00000000..412573aa --- /dev/null +++ b/app/include/cuda.hpp @@ -0,0 +1,138 @@ +#pragma once + +#include + +#define cudaErrchk(ans) \ + { cuda::Errchk((ans), __FILE__, __LINE__); } + +namespace cuda { + +inline void Errchk(cudaError_t code, const char *file, int line, bool abort = true) { + if (code != cudaSuccess) { + std::cerr + << "GPUassert: " << cudaGetErrorString(code) + << " (" << file << ":" << line << ")" + << std::endl; + if (abort) exit(code); + } +} + +template +__device__ __host__ void swap(T &a, T &b) { + T c(a); + a = b; + b = c; +}; + +template +struct pair { + U first; + V second; + __device__ __host__ bool operator<(const pair &p) const { + if (first != p.first) + return first < p.first; + else + return second < p.second; + } + pair &operator=(const std::pair &p){ + first = p.first; + second = p.second; + return *this; + } +}; + +template +class vector { + size_t capacity; + size_t sz = 0; + T *arr = nullptr; + + public: + vector(size_t cap) { + capacity = cap; + cudaErrchk(cudaMallocManaged(&arr, capacity * sizeof(T))); + } + + vector(size_t cap, size_t s, T val = T()) : vector(cap) { + sz = s; + for(size_t i = 0; i < sz; ++i) + arr[i] = val; + } + + static vector *constructShared(size_t cap){ + vector *ret; + cudaErrchk(cudaMallocManaged(&ret, sizeof(vector))); + return new (ret) vector(cap); + } + static vector *constructShared(size_t cap, size_t s, T val = T()){ + vector *ret; + cudaErrchk(cudaMallocManaged(&ret, sizeof(vector))); + return new (ret) vector(cap, s, val); + } + + void destroyShared(){ + this->~vector(); + cudaErrchk(cudaFree(this)); + } + + __device__ __host__ + T &operator[](size_t i) { + return arr[i]; + } + + __device__ __host__ + const T &operator[](size_t i) const { + return arr[i]; + } + + __device__ __host__ + T &at(size_t i) { + assert(i < sz); + return arr[i]; + } + + __device__ __host__ + const T &at(size_t i) const { + assert(i < sz); + return arr[i]; + } + + __device__ __host__ + size_t size() const { + return sz; + } + + __device__ __host__ + bool empty() const { + return sz == 0; + } + + __device__ __host__ + T &push_back(const T &val) { + assert(sz + 1 <= capacity); + return arr[sz++] = val; + } + + __device__ __host__ + void pop_back() { + assert(!empty()); + arr[--sz].~T(); + } + + template + __device__ __host__ + T &emplace_back(Args&&... args) { + assert(sz + 1 <= capacity); + return *new (arr + sz++) T(args...); + } + + __device__ __host__ + vector &operator=(const vector &v) = delete; + + ~vector(){ + while(!empty()) + arr[--sz].~T(); + cudaErrchk(cudaFree(arr)); + } +}; +} // namespace cuda diff --git a/app/include/shortest-path/DijkstraCuda.hpp b/app/include/shortest-path/DijkstraCuda.hpp new file mode 100644 index 00000000..11d459e1 --- /dev/null +++ b/app/include/shortest-path/DijkstraCuda.hpp @@ -0,0 +1,25 @@ +#pragma once + +#include + +#include "shortest-path/ShortestPathAll.hpp" + +class DijkstraCuda : public ShortestPathAll { + std::vector edges; + std::vector> adj; + + size_t numberStartNodes; + Graph::Node *startNodes; + + Graph::Edge **prev; + Graph::Edge::Weight **dist; + + public: + virtual void initialize(const Graph *G, const std::list &s); + virtual void run(); + virtual Graph::Edge getPrev(Graph::Node s, Graph::Node d) const; + virtual Graph::Edge::Weight getPathWeight(Graph::Node s, Graph::Node d) const; + virtual bool hasVisited(Graph::Node s, Graph::Node u) const; + + ~DijkstraCuda(); +}; diff --git a/app/include/shortest-path/ShortestPathAll.hpp b/app/include/shortest-path/ShortestPathAll.hpp new file mode 100644 index 00000000..c1f37bc2 --- /dev/null +++ b/app/include/shortest-path/ShortestPathAll.hpp @@ -0,0 +1,55 @@ +#pragma once + +#include + +#include "Graph.hpp" + +/** + * @brief Shortest Path From One Node to All other Nodes (Shortest Path One Many Interface) + * + */ +class ShortestPathAll { + public: + /** + * @brief Initializes the data members that are required for the algorithm's execution + * + * @param G Graph + * @param s Starting node + */ + virtual void initialize(const Graph *G, const std::list &s) = 0; + + /** + * @brief Execute the algorithm + * + */ + virtual void run() = 0; + + /** + * @brief Retrieves the node chosen prior to getting to node d + * + * @param d Destination Node + * @return Graph::Edge Edge traversed before getting to destination Node + */ + virtual Graph::Edge getPrev(Graph::Node s, Graph::Node d) const = 0; + + /** + * @brief Retrieves the sequence of nodes of the path ending at d + * + * @param d Destination Node + * @return std::list Sequence of nodes that describe the path to d + */ + virtual Graph::Path getPath(Graph::Node s, Graph::Node d) const final; + + virtual Graph::Edge::Weight getPathWeight(Graph::Node s, Graph::Node d) const = 0; + + /** + * @brief Checks if a specific node was marked as visited + * + * @param u Node to be checked + * @return true If the node has been already visited + * @return false Otherwise + */ + virtual bool hasVisited(Graph::Node s, Graph::Node u) const = 0; + + virtual ~ShortestPathAll(); +}; diff --git a/app/include/structs/BinaryHeapCuda.hpp b/app/include/structs/BinaryHeapCuda.hpp new file mode 100644 index 00000000..14d6e1e5 --- /dev/null +++ b/app/include/structs/BinaryHeapCuda.hpp @@ -0,0 +1,130 @@ +#pragma once + +#include + +#include "cuda.hpp" + +template +class BinaryHeapCuda { + template + friend class cuda::vector; + + public: + class Element { + friend BinaryHeapCuda; + template + friend class cuda::vector; + + private: + BinaryHeapCuda &binaryHeap; + size_t index; + T value; + __device__ __host__ Element(BinaryHeapCuda &heap, size_t i, T t) + : binaryHeap(heap), index(i), value(t) {} + + public: + __device__ __host__ T getValue() { return value; } + __device__ __host__ void decreaseKey(T t) { + value = t; + binaryHeap.heapifyDown(index); + } + + private: + __device__ __host__ static void swap(Element *&e1, Element *&e2) { + cuda::swap(e1->index, e2->index); + cuda::swap(e1, e2); + } + }; + + private: + typedef cuda::vector Elements; + typedef cuda::vector Container; + + Elements *elements; + Container *container; + + public: + BinaryHeapCuda(size_t s) : elements(Elements ::constructShared(s)), + container(Container::constructShared(s, 1, nullptr)) {} + + static BinaryHeapCuda *constructShared(size_t s) { + BinaryHeapCuda *ret; + cudaErrchk(cudaMallocManaged(&ret, sizeof(BinaryHeapCuda))); + return new (ret) BinaryHeapCuda(s); + } + + __device__ __host__ T top() { + return (*container)[1]->getValue(); + } + + __device__ __host__ size_t size() const { + return container->size() - 1; + } + + __device__ __host__ bool empty() const { + return size() == 0; + } + + __device__ __host__ Element &push(T t) { + Element *it = &elements->emplace_back(*this, container->size(), t); + container->push_back(it); + + heapifyDown(container->size() - 1); + + return *it; + } + + __device__ __host__ T pop() { + T ret = container->at(1)->getValue(); + + Element::swap((*container)[1], (*container)[container->size() - 1]); + // delete (*container)[container->size() - 1]; + container->pop_back(); + + heapifyUp(1); + + return ret; + } + + ~BinaryHeapCuda() { + container->destroyShared(); + } + + void destroyShared() { + this->~BinaryHeapCuda(); + cudaErrchk(cudaFree(this)); + } + + private: + __device__ __host__ void heapifyUp(size_t i) { + while (true) { + size_t l = i << 1; + size_t r = l | 1; + size_t smallest = i; + + if (l < container->size() && (*container)[l]->getValue() < (*container)[smallest]->getValue()) { + smallest = l; + } + + if (r < container->size() && (*container)[r]->getValue() < (*container)[smallest]->getValue()) { + smallest = r; + } + + if (smallest == i) break; + + Element::swap((*container)[i], (*container)[smallest]); + i = smallest; + } + } + + __device__ __host__ void heapifyDown(size_t i) { + while (i > 1) { + size_t p = i >> 1; + if ((*container)[i]->getValue() < (*container)[p]->getValue()) { + Element::swap((*container)[i], (*container)[p]); + } else + break; + i = p; + } + } +}; diff --git a/app/src/shortest-path/DijkstraCuda.cu b/app/src/shortest-path/DijkstraCuda.cu new file mode 100644 index 00000000..60ee4e14 --- /dev/null +++ b/app/src/shortest-path/DijkstraCuda.cu @@ -0,0 +1,214 @@ +#include "shortest-path/DijkstraCuda.hpp" + +#include +#include +#include +#include +#include +#include + +#include "structs/BinaryHeapCuda.hpp" + +using namespace std; + +typedef Graph::Node Node; +typedef Graph::Edge::Weight Weight; +typedef Graph::Edge Edge; +template +using umap = std::unordered_map; +typedef umap dist_t; +typedef umap prev_t; +typedef BinaryHeapCuda> MinPriorityQueue; +typedef cuda::vector Elements; +typedef std::chrono::high_resolution_clock hrc; + +void DijkstraCuda::initialize(const Graph *G, const list &s) { + const vector &nodes = G->getNodes(); + + adj.clear(); + size_t numberNodes = (nodes.empty() ? 1 : *max_element(nodes.begin(), nodes.end()) + 1); + adj.resize(numberNodes, {-1, -1}); + + edges.clear(); + size_t numberEdges = 0; + for (const Node &u : nodes) numberEdges += G->getAdj(u).size(); + edges.reserve(numberEdges); + + size_t edgeIdx = 0; + for (const Node &u : nodes) { + const auto &es = G->getAdj(u); + adj[u] = pair(edgeIdx, edgeIdx + es.size()); + edges.insert(edges.end(), es.begin(), es.end()); + edgeIdx += es.size(); + } + assert(edgeIdx == numberEdges); + + numberStartNodes = s.size(); + cudaErrchk(cudaMallocManaged(&startNodes, numberStartNodes * sizeof(Node))); + copy(s.begin(), s.end(), startNodes); + + cudaErrchk(cudaMallocManaged(&prev, numberNodes * sizeof(Edge *))); + cudaErrchk(cudaMallocManaged(&dist, numberNodes * sizeof(Weight *))); + fill(prev, prev + numberNodes, nullptr); + fill(dist, dist + numberNodes, nullptr); + for (const Node &u : s) { + cudaErrchk(cudaMallocManaged(&prev[u], numberNodes * sizeof(Edge))); + cudaErrchk(cudaMallocManaged(&dist[u], numberNodes * sizeof(Weight))); + fill(prev[u], prev[u] + numberNodes, Graph::EDGE_INVALID); + fill(dist[u], dist[u] + numberNodes, Edge::WEIGHT_INF); + } +} + +union EdgeInt4 { + int4 i; + Edge e; +}; + +union AdjInt2 { + int2 i; + cuda::pair p; +}; + +__device__ void runDijkstra( + size_t numberNodes, size_t numberEdges, + cudaTextureObject_t edges, + cudaTextureObject_t adj, + Node s, + Elements &elements, + MinPriorityQueue &Q, + Edge *prev, + Weight *dist) { + dist[s] = 0; + auto el = Q.push({0, s}); + elements[s] = ⪙ + while (!Q.empty()) { + cuda::pair p = Q.top(); Q.pop(); + Node u = p.second; + Weight du = p.first; + AdjInt2 ai2 = {.i = tex1Dfetch(adj, u)}; + for (uint32_t i = ai2.p.first; i < ai2.p.second; ++i) { + EdgeInt4 ei4 = {.i = tex1Dfetch(edges, i)}; + const Edge &e = ei4.e; + Weight c_ = du + e.w; + Weight &distV = dist[e.v]; + if (c_ < distV) { + if (elements[e.v]) + elements[e.v]->decreaseKey({c_, e.v}); + else + elements[e.v] = &Q.push({c_, e.v}); + distV = c_; + prev[e.v] = e; + } + } + } +} + +__global__ void runDijkstraKernel( + size_t numberStartNodes, Node *startNodes, + size_t numberNodes, size_t numberEdges, + cudaTextureObject_t edges, + cudaTextureObject_t adj, + cuda::vector &elements, + cuda::vector &Q, + Edge **prev, + Weight **dist) { + int a = blockIdx.x * blockDim.x + threadIdx.x; + int b = blockIdx.y * blockDim.y + threadIdx.y; + const int BMAX = blockDim.y * gridDim.y; + int i = a * BMAX + b; + if (i >= numberStartNodes) + return; + Node s = startNodes[i]; + runDijkstra(numberNodes, numberEdges, edges, adj, s, *elements.at(i), *Q.at(i), prev[s], dist[s]); +} + +void DijkstraCuda::run() { + const size_t &numberEdges = edges.size(); + const size_t &numberNodes = adj.size(); + + // Inspired by https://stackoverflow.com/q/55348493/12283316 + // Edges texture + int4 *edgesArr; + cudaTextureObject_t edgesTex; + assert(sizeof(Edge) == sizeof(int4)); + const size_t edgeSize = sizeof(int4) * numberEdges; + cudaErrchk(cudaMalloc(&edgesArr, edgeSize)); + cudaMemcpy(edgesArr, &edges[0], edgeSize, cudaMemcpyHostToDevice); + struct cudaResourceDesc edgesResDesc; + edgesResDesc.resType = cudaResourceTypeLinear; + edgesResDesc.res.linear.devPtr = edgesArr; + edgesResDesc.res.linear.sizeInBytes = edgeSize; + edgesResDesc.res.linear.desc = cudaCreateChannelDesc(); + struct cudaTextureDesc edgesTexDesc = {}; + cudaErrchk(cudaCreateTextureObject(&edgesTex, &edgesResDesc, &edgesTexDesc, NULL)); + + // Adj texture + int2 *adjArr; + cudaTextureObject_t adjTex; + assert(sizeof(cuda::pair) == sizeof(int2)); + const size_t adjSize = sizeof(int2) * numberNodes; + cudaErrchk(cudaMalloc(&adjArr, adjSize)); + cudaMemcpy(adjArr, &adj[0], adjSize, cudaMemcpyHostToDevice); + struct cudaResourceDesc adjResDesc; + adjResDesc.resType = cudaResourceTypeLinear; + adjResDesc.res.linear.devPtr = adjArr; + adjResDesc.res.linear.sizeInBytes = adjSize; + adjResDesc.res.linear.desc = cudaCreateChannelDesc(); + struct cudaTextureDesc adjTexDesc = {}; + cudaErrchk(cudaCreateTextureObject(&adjTex, &adjResDesc, &adjTexDesc, NULL)); + + // Elements + cuda::vector *elements = cuda::vector::constructShared(numberStartNodes); + for (size_t i = 0; i < numberStartNodes; ++i) + elements->emplace_back(Elements::constructShared(numberNodes, numberNodes, nullptr)); + + // Q + cuda::vector *Q = cuda::vector::constructShared(numberStartNodes); + for (size_t i = 0; i < numberStartNodes; ++i) + Q->emplace_back(MinPriorityQueue::constructShared(numberNodes)); + + const size_t &N = numberStartNodes; + dim3 threadsPerBlock(16, 8); + dim3 numBlocks( + (N + threadsPerBlock.x - 1) / threadsPerBlock.x, + (N + threadsPerBlock.y - 1) / threadsPerBlock.y); + runDijkstraKernel<<>>( + numberStartNodes, startNodes, + numberNodes, numberEdges, + edgesTex, adjTex, + *elements, *Q, + prev, dist); + cudaErrchk(cudaPeekAtLastError()); + cudaErrchk(cudaDeviceSynchronize()); + + elements->destroyShared(); + Q->destroyShared(); + + cudaErrchk(cudaFree(edgesArr)); + cudaErrchk(cudaFree(adjArr)); +} + +Edge DijkstraCuda::getPrev(Node s, Node d) const { + const size_t &numberNodes = adj.size(); + if (s >= numberNodes || prev[s] == nullptr) + throw out_of_range("s is not a valid start node"); + if (d >= numberNodes) + throw out_of_range("d is not a valid destination node"); + return prev[s][d]; +} + +Weight DijkstraCuda::getPathWeight(Node s, Node d) const { + const size_t &numberNodes = adj.size(); + if (s >= numberNodes || dist[s] == nullptr) + throw out_of_range("s is not a valid start node"); + if (d >= numberNodes) + throw out_of_range("d is not a valid destination node"); + return dist[s][d]; +} + +bool DijkstraCuda::hasVisited(Node s, Node u) const { + return getPathWeight(s, u) != Edge::WEIGHT_INF; +} + +DijkstraCuda::~DijkstraCuda() { +} diff --git a/app/src/shortest-path/ShortestPathAll.cpp b/app/src/shortest-path/ShortestPathAll.cpp new file mode 100644 index 00000000..556831cc --- /dev/null +++ b/app/src/shortest-path/ShortestPathAll.cpp @@ -0,0 +1,20 @@ +#include "shortest-path/ShortestPathAll.hpp" + +#include + +using namespace std; + +ShortestPathAll::~ShortestPathAll(){} + +Graph::Path ShortestPathAll::getPath(Graph::Node s, Graph::Node d) const{ + if(d == s) return Graph::Path(); + list res; + Graph::Edge e = getPrev(s, d); + if(e.u == Graph::NODE_INVALID) return Graph::Path({Graph::EDGE_INVALID}); + while(e.u != s){ + res.push_front(e); + e = getPrev(s, e.u); + } + res.push_front(e); + return Graph::Path(res.begin(), res.end()); +} diff --git a/app/src/static/algos/FrankWolfe.cpp b/app/src/static/algos/FrankWolfe.cpp index b5594f45..48b83f87 100644 --- a/app/src/static/algos/FrankWolfe.cpp +++ b/app/src/static/algos/FrankWolfe.cpp @@ -7,7 +7,7 @@ #include #include "convex/QuadraticSolver.hpp" -#include "shortest-path/Dijkstra.hpp" +#include "shortest-path/DijkstraCuda.hpp" using namespace std; @@ -65,29 +65,35 @@ StaticSolutionBase FrankWolfe::step1() { Graph G = problem.supply.toGraph(xn); - unordered_map> shortestPaths; + // unordered_map> shortestPaths; const vector startNodes = problem.demand.getStartNodes(); - for (const Node &u : startNodes) { - shortestPaths.emplace(u, new Dijkstra()); - shortestPaths[u].get()->initialize(&G, u); - } + // for (const Node &u : startNodes) { + // shortestPaths.emplace(u, new Dijkstra()); + // shortestPaths[u].get()->initialize(&G, u); + // } + + // vector> results; + // for (const Node &u : startNodes) { + // results.emplace_back(pool.push([&shortestPaths, u](int) { + // ShortestPathOneMany *sp = shortestPaths.at(u).get(); + // sp->run(); + // })); + // } + // for (future &r : results) r.get( + + list startNodesGraph; + for(const Node &u: startNodes) + startNodesGraph.push_back((Graph::Node)u); + + ShortestPathAll *sp = new DijkstraCuda(); + sp->initialize(&G, startNodesGraph); + sp->run(); - vector> results; for (const Node &u : startNodes) { - results.emplace_back(pool.push([&shortestPaths, u](int) { - ShortestPathOneMany *sp = shortestPaths.at(u).get(); - sp->run(); - })); - } - for (future &r : results) r.get(); - - for (const Node &u : startNodes) { - const ShortestPathOneMany *sp = shortestPaths[u].get(); - const vector endNodes = problem.demand.getDestinations(u); for (const Node &v : endNodes) { - Graph::Path path = sp->getPath(v); + Graph::Path path = sp->getPath((Graph::Node)u, (Graph::Node)v); if (path.size() == 1 && path.front().id == Graph::EDGE_INVALID.id) throw logic_error("Could not find path " + to_string(u) + "->" + to_string(v)); @@ -102,6 +108,8 @@ StaticSolutionBase FrankWolfe::step1() { } } + delete sp; + return xstar; } diff --git a/app/test/test_Dijkstra.cpp b/app/test/test_Dijkstra.cpp index a1916254..b04a8a7c 100644 --- a/app/test/test_Dijkstra.cpp +++ b/app/test/test_Dijkstra.cpp @@ -1,9 +1,10 @@ -#include #include +#include #include #include "data/sumo/TAZs.hpp" #include "shortest-path/Dijkstra.hpp" +#include "shortest-path/DijkstraCuda.hpp" #include "static/algos/AllOrNothing.hpp" #include "static/supply/BPRNetwork.hpp" @@ -48,10 +49,9 @@ void testPath(std::vector expected, Graph::Path got) { } TEST_CASE("Dijkstra's algorithm", "[shortestpath][shortestpath-onemany][dijkstra]") { - SECTION("Start 0") { Graph G = graph1(); - + ShortestPathOneMany *shortestPath = new Dijkstra(); shortestPath->initialize(&G, 0); shortestPath->run(); @@ -117,25 +117,25 @@ TEST_CASE("Dijkstra's algorithm", "[shortestpath][shortestpath-onemany][dijkstra sp.get()->run(); const double v1 = 13.89, l1 = 14.07; - const double v2 = 8.33, l2 = 18.80; + const double v2 = 8.33, l2 = 18.80; const double v3 = 13.89, l3 = 33.24; - const double v4 = 8.33, l4 = 39.34; - const double t1 = l1/(v1*0.9); - const double t2 = l2/(v2*0.9); - const double t3 = l3/(v3*0.9); - const double t4 = l4/(v4*0.9); + const double v4 = 8.33, l4 = 39.34; + const double t1 = l1 / (v1 * 0.9); + const double t2 = l2 / (v2 * 0.9); + const double t3 = l3 / (v3 * 0.9); + const double t4 = l4 / (v4 * 0.9); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("2").first), WithinAbs(0, 1e-6)); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("2").second), WithinAbs(t2, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").first), WithinAbs(t2+10, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").second), WithinAbs(t2+10+t1, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").first), WithinAbs(t2 + 10, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").second), WithinAbs(t2 + 10 + t1, 1e-6)); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").first), WithinAbs(t2, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").second), WithinAbs(t2+t4, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").second), WithinAbs(t2 + t4, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").first), WithinAbs(t2+20, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").second), WithinAbs(t2+20+t3, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").first), WithinAbs(t2 + 20, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").second), WithinAbs(t2 + 20 + t3, 1e-6)); delete network; } @@ -157,34 +157,34 @@ TEST_CASE("Dijkstra's algorithm", "[shortestpath][shortestpath-onemany][dijkstra sp.get()->run(); const double v1 = 13.89, l1 = 14.07; - const double v2 = 8.33, l2 = 18.80; + const double v2 = 8.33, l2 = 18.80; const double v3 = 13.89, l3 = 33.24; - const double v4 = 8.33, l4 = 39.34; - const double t1 = l1/(v1*0.9); - const double t2 = l2/(v2*0.9); - const double t3 = l3/(v3*0.9); - const double t4 = l4/(v4*0.9); + const double v4 = 8.33, l4 = 39.34; + const double t1 = l1 / (v1 * 0.9); + const double t2 = l2 / (v2 * 0.9); + const double t3 = l3 / (v3 * 0.9); + const double t4 = l4 / (v4 * 0.9); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("2").first), WithinAbs(0, 1e-6)); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("2").second), WithinAbs(t2, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").first), WithinAbs(t2+10, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").second), WithinAbs(t2+10+t1, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").first), WithinAbs(t2 + 10, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-1").second), WithinAbs(t2 + 10 + t1, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("1").first), WithinAbs(t2+10+t1+20, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("1").second), WithinAbs(t2+10+t1+20+t1, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("1").first), WithinAbs(t2 + 10 + t1 + 20, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("1").second), WithinAbs(t2 + 10 + t1 + 20 + t1, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").first), WithinAbs(t2+10+t1+20+t1, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").second), WithinAbs(t2+10+t1+20+t1+t3, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").first), WithinAbs(t2 + 10 + t1 + 20 + t1, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-3").second), WithinAbs(t2 + 10 + t1 + 20 + t1 + t3, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("3").first), WithinAbs(t2+10+t1+20+t1+t3+20, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("3").second), WithinAbs(t2+10+t1+20+t1+t3+20+t3, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("3").first), WithinAbs(t2 + 10 + t1 + 20 + t1 + t3 + 20, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("3").second), WithinAbs(t2 + 10 + t1 + 20 + t1 + t3 + 20 + t3, 1e-6)); REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").first), WithinAbs(t2, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").second), WithinAbs(t2+t4, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("-4").second), WithinAbs(t2 + t4, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("4").first), WithinAbs(t2+t4+20, 1e-6)); - REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("4").second), WithinAbs(t2+t4+20+t4, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("4").first), WithinAbs(t2 + t4 + 20, 1e-6)); + REQUIRE_THAT(sp.get()->getPathWeight(adapter.toNodes("4").second), WithinAbs(t2 + t4 + 20 + t4, 1e-6)); delete network; } @@ -212,3 +212,56 @@ TEST_CASE("Dijkstra's algorithm", "[shortestpath][shortestpath-onemany][dijkstra delete network; } } + +TEST_CASE("Dijkstra's algorithm (CUDA)", "[dijkstra-cuda]") { + SECTION("Start 0") { + Graph G = graph1(); + + ShortestPathAll *shortestPath = new DijkstraCuda(); + shortestPath->initialize(&G, {0}); + shortestPath->run(); + + testPath({0}, shortestPath->getPath(0, 0)); + testPath({0, 1}, shortestPath->getPath(0, 1)); + testPath({0, 1, 2}, shortestPath->getPath(0, 2)); + testPath({0, 1, 2, 3}, shortestPath->getPath(0, 3)); + testPath({0, 1, 2, 3, 4}, shortestPath->getPath(0, 4)); + testPath({0, 1, 2, 5}, shortestPath->getPath(0, 5)); + testPath({0, 1, 2, 5, 6}, shortestPath->getPath(0, 6)); + + REQUIRE_THAT(shortestPath->getPathWeight(0, 0), WithinAbs(0, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 1), WithinAbs(1, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 2), WithinAbs(3, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 3), WithinAbs(4, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 4), WithinAbs(6, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 5), WithinAbs(5, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(0, 6), WithinAbs(9, 1e-10)); + + delete shortestPath; + } + SECTION("Start 1") { + Graph G = graph1(); + + ShortestPathAll *shortestPath = new DijkstraCuda(); + shortestPath->initialize(&G, {1}); + shortestPath->run(); + + testPath({}, shortestPath->getPath(1, 0)); + testPath({1}, shortestPath->getPath(1, 1)); + testPath({1, 2}, shortestPath->getPath(1, 2)); + testPath({1, 2, 3}, shortestPath->getPath(1, 3)); + testPath({1, 2, 3, 4}, shortestPath->getPath(1, 4)); + testPath({1, 2, 5}, shortestPath->getPath(1, 5)); + testPath({1, 2, 5, 6}, shortestPath->getPath(1, 6)); + + REQUIRE_THAT(shortestPath->getPathWeight(1, 0), WithinAbs(Graph::Edge::WEIGHT_INF, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 1), WithinAbs(0, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 2), WithinAbs(2, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 3), WithinAbs(3, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 4), WithinAbs(5, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 5), WithinAbs(4, 1e-10)); + REQUIRE_THAT(shortestPath->getPathWeight(1, 6), WithinAbs(8, 1e-10)); + + delete shortestPath; + } +}