GSParLib source and examples

.github/workflows/gspar-build-ci.yml

@@ -0,0 +1,26 @@
+name: GSPar Build CI
+
+on:
+  push:
+    branches: [ main ]
+  pull_request:
+    branches: [ main ]
+
+jobs:
+  build:
+
+    runs-on: ubuntu-latest
+
+    steps:
+    - name: Check out repository
+      uses: actions/checkout@v2
+    - name: Fix apt on GitHub Actions
+      run: sudo gem install apt-spy2 && sudo apt-spy2 fix --commit --launchpad --country=US
+    - name: Update apt
+      run: sudo apt-get update
+    - name: Install OpenCL
+      run: sudo apt-get -o Acquire::Retries=3 install opencl-headers nvidia-opencl-dev #nvidia-libopencl1-384
+    - name: Install CUDA
+      run: sudo apt-get -o Acquire::Retries=3 install nvidia-cuda-dev
+    - name: Build library
+      run: make

.gitignore

@@ -0,0 +1,40 @@
+# VS Code
+.vscode/ipch/*
+
+# Prerequisites
+*.d
+
+# Compiled Object files
+*.slo
+*.lo
+*.o
+*.obj
+
+# Precompiled Headers
+*.gch
+*.pch
+
+# Compiled Dynamic libraries
+*.so
+*.dylib
+*.dll
+
+# Fortran module files
+*.mod
+*.smod
+
+# Compiled Static libraries
+*.lai
+*.la
+*.a
+*.lib
+
+# Ignore the build and lib dirs
+build
+lib/*
+
+# Executables
+bin/*
+
+# Auto-generated
+generated_*.cpp

LICENSE

@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (c) 2023 Parallel Applications Modelling Group - GMAP
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

Makefile

@@ -0,0 +1,151 @@
+# Compilers
+COMPILER := g++
+# Directories
+SRCDIR := src
+BUILDDIR := build
+TARGETDIR := bin
+EXAMPLESDIR := examples
+EXAMPLEDRIVERAPIDIR := $(EXAMPLESDIR)/driver_api
+EXAMPLEPATTERNAPIDIR := $(EXAMPLESDIR)/pattern_api
+EXAMPLESEQUENTIALDIR := $(EXAMPLESDIR)/sequential
+THIRDPTDIR := thirdpt
+MARX2DIR := $(THIRDPTDIR)/marX2
+LIBMARX2PATH := $(MARX2DIR)/libmarX2.a
+# Names
+LIBNAME := gspar
+GSPARNAME := gspar
+CUDANAME := cuda
+OCLNAME := opencl
+DRIVERAPINAME := driverapi
+PATTERNAPINAME := patternapi
+SEQUENTIALNAME := seq
+# App names
+MANDELNAME := mandel
+LANEDETECTIONNAME := lanedetection
+# Target
+TARGET := $(TARGETDIR)/lib$(LIBNAME).so
+# Others
+SPACE := 
+SPACE +=
+SRCEXT := cpp
+EXAMPLESTARGETPREFIX := ex
+
+# Files
+SOURCES := $(wildcard $(SRCDIR)/*.$(SRCEXT))
+OBJECTS := $(patsubst $(SRCDIR)/%,$(BUILDDIR)/%,$(SOURCES:.$(SRCEXT)=.o))
+
+CFLAGS := -Wall -std=c++14 -O3 -Wno-reorder
+LIB := -Llib -L/usr/local/cuda/targets/x86_64-linux/lib -L/usr/local/cuda/targets/x86_64-linux/lib/stubs -L/usr/local/cuda/lib -L/usr/local/cuda/lib64
+LIBOCL := -lOpenCL
+LIBCUDADRIVER := -lcuda
+LIBCUDANVRTC := -lnvrtc
+LIBPTHREAD := -pthread
+PATHSLIB := -I/usr/local/cuda/include -Isrc
+PATHSTEST := $(PATHSLIB) -I$(THIRDPTDIR) -I$(EXAMPLESDIR)/include
+TESTLIB := -L$(TARGETDIR) -l$(LIBNAME)
+
+PATHOPENCV := -I/usr/local/include/opencv4
+LIBSOPENCV := -lopencv_core -lopencv_highgui -lopencv_imgproc -lopencv_imgcodecs
+EXTRADEPS := 
+INCMANDEL := 
+LIBMANDEL := 
+ifdef DEBUG
+	DEFS +=-DDEBUG -DGSPAR_DEBUG
+	EXTRADEPS := $(LIBMARX2PATH)
+	INCMANDEL := -I$(MARX2DIR) -L$(MARX2DIR)
+	LIBMANDEL := -lmarX2 -lX11 -lm
+endif
+
+CLR_BLUE := \033[0;34m
+CLR_ORANGE := \033[0;33m
+CLR_DARKCYAN := \033[0;36m
+CLR_NO := \033[0m
+
+# Functions
+get_paths_mandel = $(if $(findstring $(MANDELNAME), $(1)), $(INCMANDEL))
+get_paths_lanedetection = $(if $(findstring $(LANEDETECTIONNAME), $(1)), $(PATHOPENCV))
+get_paths = $(strip $(PATHSTEST) $(LIB) $(call get_paths_mandel, $(1)) $(call get_paths_lanedetection, $(1)) )
+
+get_libs_mandel = $(if $(findstring $(MANDELNAME), $(1)), $(LIBMANDEL))
+get_libs_lanedetection = $(if $(findstring $(LANEDETECTIONNAME), $(1)), $(LIBSOPENCV))
+get_libs = $(strip $(call get_libs_mandel, $(1)) $(call get_libs_lanedetection, $(1)))
+
+# Driver API examples
+EXAMPLESOURCES_DRIVERAPI := $(wildcard $(EXAMPLEDRIVERAPIDIR)/*.$(SRCEXT))
+EXAMPLETARGETS_DRIVERAPI_CUDA := $(patsubst $(EXAMPLEDRIVERAPIDIR)/%,$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%,$(EXAMPLESOURCES_DRIVERAPI:.$(SRCEXT)=_$(CUDANAME)))
+EXAMPLETARGETS_DRIVERAPI_OPENCL := $(patsubst $(EXAMPLEDRIVERAPIDIR)/%,$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%,$(EXAMPLESOURCES_DRIVERAPI:.$(SRCEXT)=_$(OCLNAME)))
+# Pattern API examples
+EXAMPLESOURCES_PATTERNAPI := $(wildcard $(EXAMPLEPATTERNAPIDIR)/*.$(SRCEXT))
+EXAMPLETARGETS_PATTERNAPI_CUDA := $(patsubst $(EXAMPLEPATTERNAPIDIR)/%,$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%,$(EXAMPLESOURCES_PATTERNAPI:.$(SRCEXT)=_$(CUDANAME)))
+EXAMPLETARGETS_PATTERNAPI_OPENCL := $(patsubst $(EXAMPLEPATTERNAPIDIR)/%,$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%,$(EXAMPLESOURCES_PATTERNAPI:.$(SRCEXT)=_$(OCLNAME)))
+# Sequential examples
+EXAMPLESOURCES_SEQUENTIAL := $(wildcard $(EXAMPLESEQUENTIALDIR)/*.$(SRCEXT))
+EXAMPLETARGETS_SEQUENTIAL := $(patsubst $(EXAMPLESEQUENTIALDIR)/%,$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(SEQUENTIALNAME)_%,$(EXAMPLESOURCES_SEQUENTIAL:.$(SRCEXT)=))
+
+
+# Build targets
+
+$(TARGET): $(OBJECTS) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Linking dynamic library ${CLR_ORANGE}$(TARGET)${CLR_NO}..."
+	$(COMPILER) $(DEFS) -shared -fPIC -o $(TARGET) $^ $(LIB) $(LIBOCL) $(LIBCUDADRIVER) $(LIBCUDANVRTC)
+
+$(BUILDDIR)/%.o: $(SRCDIR)/%.$(SRCEXT) | $(BUILDDIR)
+	@echo "${CLR_DARKCYAN}Compiling and assembling object ${CLR_ORANGE}$@${CLR_NO}..."
+	$(COMPILER) $(DEFS) $(CFLAGS) $(PATHSLIB) -c -fPIC -o $@ $<
+
+$(TARGETDIR):
+	@mkdir -p $@
+
+$(BUILDDIR):
+	@mkdir -p $@
+
+
+.PHONY: examples
+examples: examples_driver_api examples_pattern_api examples_sequential
+
+# Driver API examples
+examples_driver_api: $(EXAMPLETARGETS_DRIVERAPI_CUDA) $(EXAMPLETARGETS_DRIVERAPI_OPENCL)
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%: $(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%_$(CUDANAME) $(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%_$(OCLNAME) ;
+# Lib to CUDA
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%_$(CUDANAME): $(EXAMPLEDRIVERAPIDIR)/%.$(SRCEXT) $(TARGET) $(EXTRADEPS) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Building GSPar Driver API example ${CLR_ORANGE}$@${CLR_DARKCYAN} from $<${CLR_NO}"
+	$(COMPILER) $(DEFS) -DGSPARDRIVER_CUDA $(CFLAGS) $< $(call get_paths, $<) $(TESTLIB) -o $@ $(LIBPTHREAD) $(call get_libs, $<)
+# Lib to OpenCL
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(DRIVERAPINAME)_%_$(OCLNAME): $(EXAMPLEDRIVERAPIDIR)/%.$(SRCEXT) $(TARGET) $(EXTRADEPS) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Building GSPar Driver API example ${CLR_ORANGE}$@${CLR_DARKCYAN} from $<${CLR_NO}"
+	$(COMPILER) $(DEFS) -DGSPARDRIVER_OPENCL $(CFLAGS) $< $(call get_paths, $<) $(TESTLIB) -o $@ $(LIBPTHREAD) $(call get_libs, $<)
+
+# Pattern API examples
+examples_pattern_api: $(EXAMPLETARGETS_PATTERNAPI_CUDA) $(EXAMPLETARGETS_PATTERNAPI_OPENCL)
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%: $(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%_$(CUDANAME) $(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%_$(OCLNAME) ;
+# Lib to CUDA
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%_$(CUDANAME): $(EXAMPLEPATTERNAPIDIR)/%.$(SRCEXT) $(TARGET) $(EXTRADEPS) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Building GSPar Pattern API example ${CLR_ORANGE}$@${CLR_DARKCYAN} from $<${CLR_NO}"
+	$(COMPILER) $(DEFS) -DGSPARDRIVER_CUDA $(CFLAGS) $< $(call get_paths, $<) $(TESTLIB) -o $@ $(LIBPTHREAD) $(call get_libs, $<)
+# Lib to OpenCL
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(PATTERNAPINAME)_%_$(OCLNAME): $(EXAMPLEPATTERNAPIDIR)/%.$(SRCEXT) $(TARGET) $(EXTRADEPS) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Building GSPar Pattern API example ${CLR_ORANGE}$@${CLR_DARKCYAN} from $<${CLR_NO}"
+	$(COMPILER) $(DEFS) -DGSPARDRIVER_OPENCL $(CFLAGS) $< $(call get_paths, $<) $(TESTLIB) -o $@ $(LIBPTHREAD) $(call get_libs, $<)
+
+# Sequential examples
+examples_sequential: $(EXAMPLETARGETS_SEQUENTIAL)
+$(TARGETDIR)/$(EXAMPLESTARGETPREFIX)_$(SEQUENTIALNAME)_%: $(EXAMPLESEQUENTIALDIR)/%.$(SRCEXT) | $(TARGETDIR)
+	@echo "${CLR_DARKCYAN}Building sequential example ${CLR_ORANGE}$@${CLR_DARKCYAN} from $<${CLR_NO}"
+	$(COMPILER) $(DEFS) $(CFLAGS) $< $(call get_paths, $<) -o $@ $(call get_libs, $<)
+
+$(LIBMARX2PATH): $(MARX2DIR)/marX2.c $(MARX2DIR)/marX2.h
+	@echo "${CLR_DARKCYAN}Building ${CLR_ORANGE}$(LIBMARX2PATH)${CLR_DARKCYAN}${CLR_NO}"
+	gcc -c -Wall -O3 -I/usr/X11R6/include -I$(MARX2DIR) $(MARX2DIR)/marX2.c -o $(MARX2DIR)/marX2.o
+	ar -rv $(LIBMARX2PATH) $(MARX2DIR)/marX2.o
+	ranlib $(LIBMARX2PATH)
+
+.PHONY: clean
+clean:
+	@echo "${CLR_DARKCYAN}Cleaning...${CLR_NO}"; 
+	$(RM) -r $(BUILDDIR) $(TARGETDIR) $(MARX2DIR)/*.a $(MARX2DIR)/*.o
+
+clean_lib:
+	@echo "${CLR_DARKCYAN}Cleaning lib $(TARGET)...${CLR_NO}"; 
+	$(RM) $(OBJECTS) $(TARGET)
+
+all: $(TARGET) examples

README.md

@@ -1 +1,23 @@
-# GSParLib
+# GSParLib
+
+GSParLib is a C++ object-oriented multi-level API for GPU programming that allows code portability between different GPU platforms and targets stream and data parallelism.
+
+The scientific article presenting GSParLib is currently under review.
+
+## Compilation
+
+- `make` builds the library
+- `make examples` builds the examples for both Pattern and Driver API, as well as the sequential versions. To compile just a specific set of examples, use one of:
+  - `make examples_driver_api`
+  - `make examples_pattern_api`
+  - `make examples_sequential`
+- Alternatively, it is possible to compile individual examples by referring directly to their compiled names (the `cuda`/`opencl` suffix may be ommited). Ex.: `make bin/ex_driverapi_gpuinfo` compiles both CUDA and OpenCL versions of the [gpuinfo.cpp](examples/driver_api/gpuinfo.cpp) example.
+
+To compile with debugging enabled, use `DEBUG=1 make` (both when compiling the library and the examples). This automatically enables debugging flags, so that GSParLib print various debugging information during execution.
+
+## Run examples
+
+After building the library it is necessary to make it available at runtime.
+To do this, execute `export LD_LIBRARY_PATH=<path>/bin:$LD_LIBRARY_PATH`, replacing `<path>` with the path to the repository's root folder.
+
+After this, just execute any example under the path `bin/ex_`

examples/driver_api/atomic_operations.cpp

@@ -0,0 +1,113 @@
+#include <iostream>
+#include <chrono>
+
+#ifdef GSPARDRIVER_OPENCL
+    #include "GSPar_OpenCL.hpp"
+    using namespace GSPar::Driver::OpenCL;
+#else
+    #include "GSPar_CUDA.hpp"
+    using namespace GSPar::Driver::CUDA;
+#endif
+
+const char* kernelSource = GSPAR_STRINGIZE_SOURCE(
+    GSPAR_DEVICE_KERNEL void atomicops_kernel(const int max,
+            GSPAR_DEVICE_GLOBAL_MEMORY const int *vector,
+            GSPAR_DEVICE_GLOBAL_MEMORY int *result) {
+        size_t gid = gspar_get_global_id(0);
+        if (gid <= max) {
+            gspar_atomic_add_int(result, vector[gid]);
+        }
+    }
+);
+
+void print_vector(int size, const int* vector, bool compact = false) {
+    if (compact || size > 100) {
+        std::cout << vector[0] << "..." << vector[size-1];
+    } else {
+        for (int i = 0; i < size; i++) {
+            std::cout << vector[i] << " ";
+        }
+    }
+    std::cout << std::endl;
+}
+
+int main(int argc, const char * argv[]) {
+
+    std::cout << "Testing atomic operations in GSParLib Driver API" << std::endl;
+
+    const int VECTOR_SIZE = 20;
+
+    // Create memory objects
+    int correctResult = 0;
+    int* result = new int;
+    int* vector = new int[VECTOR_SIZE];
+    for (int i = 0; i < VECTOR_SIZE; i++) {
+        vector[i] = (int)i;
+        correctResult += i;
+    }
+
+    std::cout << "Vector with " << VECTOR_SIZE << " elements:" << std::endl;
+    print_vector(VECTOR_SIZE, vector);
+
+    try {
+
+        auto t_start = std::chrono::steady_clock::now();
+
+        auto driver = Instance::getInstance();
+        driver->init();
+
+        int numGpus = driver->getGpuCount();
+        if (numGpus == 0) {
+            std::cout << "No GPU found, interrupting test" << std::endl;
+            exit(-1);
+        }
+
+        // Get the first GPU
+        auto gpu = driver->getGpu(0);
+
+        auto vector_dev = gpu->malloc(sizeof(int) * VECTOR_SIZE, vector);
+        // Async copy
+        // vector_dev->copyInAsync();
+        // vector_dev->waitAsync();
+        // Sync copy
+        vector_dev->copyIn();
+
+        auto result_dev = gpu->malloc(sizeof(int), result);
+
+        auto kernel = gpu->prepareKernel(kernelSource, "atomicops_kernel");
+        // auto kernel = new Kernel(gpu, kernelSource, "atomicops_kernel");
+
+        // Set a fixed number of threads per block for the X dimension
+        // kernel->setNumThreadsPerBlockForX(5);
+        kernel->setParameter(sizeof(VECTOR_SIZE), &VECTOR_SIZE);
+        kernel->setParameter(vector_dev);
+        kernel->setParameter(result_dev);
+
+        kernel->runAsync({VECTOR_SIZE, 0});
+        kernel->waitAsync();
+
+        result_dev->copyOut();
+
+        delete kernel;
+        delete vector_dev;
+        delete result_dev;
+
+        auto t_end = std::chrono::steady_clock::now();
+
+        // Output the result buffer
+        std::cout << "Expected result: " << correctResult << std::endl;
+        std::cout << "Actual result:   " << *result << std::endl;
+
+        delete vector;
+        delete result;
+
+        std::cout << "Test finished succesfully in " << std::chrono::duration_cast<std::chrono::milliseconds>(t_end - t_start).count() << " ms " << std::endl;
+
+        return 0;
+
+    } catch (GSPar::GSParException &ex) {
+        std::cerr << "Exception: " << ex.what() << " - " << ex.getDetails() << std::endl;
+        exit(-1);
+    }
+
+}