CUDA

.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

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

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 $<$<COMPILE_LANGUAGE:CUDA>:
+                       -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 $<$<COMPILE_LANGUAGE:CUDA>:
+                       -g
+                       -G
+                       >)
 
 # ################ SCRIPT ################
 find_package(nlohmann_json 3.2.0 REQUIRED)

app/include/Graph.hpp

@@ -2,17 +2,19 @@
 
 #include <unordered_map>
 #include <vector>
+#include <cstdint>
+#include <cstdint>
 
 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;
 

app/include/cuda.hpp

@@ -0,0 +1,138 @@
+#pragma once
+
+#include <exception>
+
+#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 <class T>
+__device__ __host__ void swap(T &a, T &b) {
+    T c(a);
+    a = b;
+    b = c;
+};
+
+template <class U, class V>
+struct pair {
+    U first;
+    V second;
+    __device__ __host__ bool operator<(const pair<U, V> &p) const {
+        if (first != p.first)
+            return first < p.first;
+        else
+            return second < p.second;
+    }
+    pair<U, V> &operator=(const std::pair<U, V> &p){
+        first = p.first;
+        second = p.second;
+        return *this;
+    }
+};
+
+template <class T>
+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<T> *constructShared(size_t cap){
+        vector<T> *ret;
+        cudaErrchk(cudaMallocManaged(&ret, sizeof(vector<T>)));
+        return new (ret) vector<T>(cap);
+    }
+    static vector<T> *constructShared(size_t cap, size_t s, T val = T()){
+        vector<T> *ret;
+        cudaErrchk(cudaMallocManaged(&ret, sizeof(vector<T>)));
+        return new (ret) vector<T>(cap, s, val);
+    }
+
+    void destroyShared(){
+        this->~vector<T>();
+        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<class... Args>
+    __device__ __host__
+    T &emplace_back(Args&&... args) {
+        assert(sz + 1 <= capacity);
+        return *new (arr + sz++) T(args...);
+    }
+
+    __device__ __host__
+    vector<T> &operator=(const vector<T> &v) = delete;
+
+    ~vector(){
+        while(!empty())
+            arr[--sz].~T();
+        cudaErrchk(cudaFree(arr));
+    }
+};
+}  // namespace cuda

app/include/shortest-path/DijkstraCuda.hpp

@@ -0,0 +1,25 @@
+#pragma once
+
+#include <utility>
+
+#include "shortest-path/ShortestPathAll.hpp"
+
+class DijkstraCuda : public ShortestPathAll {
+    std::vector<Graph::Edge> edges;
+    std::vector<std::pair<uint32_t, uint32_t>> adj;
+
+    size_t numberStartNodes;
+    Graph::Node *startNodes;
+
+    Graph::Edge **prev;
+    Graph::Edge::Weight **dist;
+
+   public:
+    virtual void initialize(const Graph *G, const std::list<Graph::Node> &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();
+};

app/include/shortest-path/ShortestPathAll.hpp

@@ -0,0 +1,55 @@
+#pragma once
+
+#include <list>
+
+#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<Graph::Node> &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<Graph::Node> 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();
+};