From 00ec9a7db62b83b4fec59e32f113901c4a72b985 Mon Sep 17 00:00:00 2001 From: svkurowski Date: Thu, 27 Nov 2014 10:15:20 +0100 Subject: [PATCH 1/2] Integrate CUDA into buildsystem and add example function Former-commit-id: 392acb148a0033cf3cb24ed9f558a7c9b625935b --- CMakeLists.txt | 45 +++++++++++++++++++++++++++++++++ cuda/kernels/graph.cu | 58 +++++++++++++++++++++++++++++++++++++++++++ cuda/kernels/graph.h | 6 +++++ storm-config.h.in | 3 +++ 4 files changed, 112 insertions(+) create mode 100644 cuda/kernels/graph.cu create mode 100644 cuda/kernels/graph.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 4c2d132b8..73c025b88 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -73,6 +73,12 @@ else() set(ENABLE_GUROBI ON) endif() +if ("${CUDA_ROOT}" STREQUAL "") + set(ENABLE_CUDA OFF) +else() + set(ENABLE_CUDA ON) +endif() + if ("${Z3_ROOT}" STREQUAL "") set(ENABLE_Z3 OFF) else() @@ -195,6 +201,13 @@ else() set(STORM_CPP_GUROBI_DEF "undef") endif() +# CUDA Defines +if (ENABLE_CUDA) + set(STORM_CPP_CUDA_DEF "define") +else() + set(STORM_CPP_CUDA_DEF "undef") +endif() + # glpk defines set(STORM_CPP_GLPK_DEF "define") @@ -227,6 +240,28 @@ configure_file ( # Add the binary dir include directory for storm-config.h include_directories("${PROJECT_BINARY_DIR}/include") +############################################################# +## +## CUDA Library generation +## +############################################################# + +if(ENABLE_CUDA) + find_package(CUDA REQUIRED) + + set(STORM_CUDA_LIB_NAME "storm-cuda") + + file(GLOB_RECURSE STORM_CUDA_KERNEL_FILES ${PROJECT_SOURCE_DIR}/cuda/kernels/*.cu) + + set(CUDA_PROPAGATE_HOST_FLAGS OFF) + set(CUDA_NVCC_FLAGS "-arch=sm_13") + + cuda_add_library(${STORM_CUDA_LIB_NAME} + ${STORM_CUDA_KERNEL_FILES} + SHARED + ) +endif() + ############################################################# ## ## Source file aggregation and clustering @@ -387,6 +422,16 @@ if (ENABLE_GUROBI) target_link_libraries(storm-performance-tests "gurobi56") endif(ENABLE_GUROBI) +############################################################# +## +## CUDA (optional) +## +############################################################# +if (ENABLE_CUDA) + message (STATUS "StoRM - Linking with CUDA") + target_link_libraries(storm ${STORM_CUDA_LIB_NAME}) +endif(ENABLE_CUDA) + ############################################################# ## ## glpk diff --git a/cuda/kernels/graph.cu b/cuda/kernels/graph.cu new file mode 100644 index 000000000..c3d708b14 --- /dev/null +++ b/cuda/kernels/graph.cu @@ -0,0 +1,58 @@ +#include "cuda/kernels/graph.h" + +#include + +const int N = 16; +const int blocksize = 16; + +#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } +inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) +{ + if (code != cudaSuccess) + { + fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) exit(code); + } +} + +__global__ +void hello(char *a, int *b) +{ + a[threadIdx.x] += b[threadIdx.x]; +} + +void helloWorldCuda() +{ + printf("CUDA TEST START\n"); + printf("Should print \"Hello World\"\n"); + + char a[N] = "Hello \0\0\0\0\0\0"; + int b[N] = {15, 10, 6, 0, -11, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + char c[N] = "YELLO \0\0\0\0\0\0"; + + char *ad; + int *bd; + const int csize = N*sizeof(char); + const int isize = N*sizeof(int); + + printf("%s", a); + + cudaMalloc( (void**)&ad, csize ); + cudaMalloc( (void**)&bd, isize ); + cudaMemcpy( ad, a, csize, cudaMemcpyHostToDevice ); + cudaMemcpy( bd, b, isize, cudaMemcpyHostToDevice ); + + dim3 dimBlock( blocksize, 1 ); + dim3 dimGrid( 1, 1 ); + hello<<>>(ad, bd); + + gpuErrchk( cudaPeekAtLastError() ); + gpuErrchk( cudaDeviceSynchronize() ); + + cudaMemcpy( c, ad, csize, cudaMemcpyDeviceToHost ); + cudaFree( ad ); + cudaFree( bd ); + + printf("%s\n", c); + printf("CUDA TEST END\n"); +} diff --git a/cuda/kernels/graph.h b/cuda/kernels/graph.h new file mode 100644 index 000000000..8f9cf8e8b --- /dev/null +++ b/cuda/kernels/graph.h @@ -0,0 +1,6 @@ +#ifndef CUDA_KERNELS_GRAPH_H +#define CUDA_KERNELS_GRAPH_H + +void helloWorldCuda(); + +#endif /* CUDA_KERNELS_GRAPH_H */ \ No newline at end of file diff --git a/storm-config.h.in b/storm-config.h.in index cb90e5af6..40db93554 100644 --- a/storm-config.h.in +++ b/storm-config.h.in @@ -17,6 +17,9 @@ // Whether Gurobi is available and to be used (define/undef) #@STORM_CPP_GUROBI_DEF@ STORM_HAVE_GUROBI +// Whether CUDA is available (define/undef) +#@STORM_CPP_CUDA_DEF@ STORM_HAVE_CUDA + // Whether GLPK is available and to be used (define/undef) #@STORM_CPP_GLPK_DEF@ STORM_HAVE_GLPK From a4a15dd7748fde2272c4d4a1cfbdf5da1ea96757 Mon Sep 17 00:00:00 2001 From: svkurowski Date: Thu, 27 Nov 2014 10:15:30 +0100 Subject: [PATCH 2/2] Add general setting to enable CUDA on runtime Former-commit-id: fc2c8e8d0c9f2b0816e269fd07c288c8765beb49 --- src/settings/modules/GeneralSettings.cpp | 6 ++++++ src/settings/modules/GeneralSettings.h | 8 ++++++++ 2 files changed, 14 insertions(+) diff --git a/src/settings/modules/GeneralSettings.cpp b/src/settings/modules/GeneralSettings.cpp index d04ec0d25..6ec11df68 100644 --- a/src/settings/modules/GeneralSettings.cpp +++ b/src/settings/modules/GeneralSettings.cpp @@ -45,6 +45,7 @@ namespace storm { const std::string GeneralSettings::statisticsOptionShortName = "stats"; const std::string GeneralSettings::bisimulationOptionName = "bisimulation"; const std::string GeneralSettings::bisimulationOptionShortName = "bisim"; + const std::string GeneralSettings::cudaOptionName = "cuda"; GeneralSettings::GeneralSettings(storm::settings::SettingsManager& settingsManager) : ModuleSettings(settingsManager, moduleName) { this->addOption(storm::settings::OptionBuilder(moduleName, helpOptionName, false, "Shows all available options, arguments and descriptions.").setShortName(helpOptionShortName) @@ -95,6 +96,7 @@ namespace storm { this->addOption(storm::settings::OptionBuilder(moduleName, constantsOptionName, false, "Specifies the constant replacements to use in symbolic models. Note that Note that this requires the model to be given as an symbolic model (i.e., via --" + symbolicOptionName + ").").setShortName(constantsOptionShortName) .addArgument(storm::settings::ArgumentBuilder::createStringArgument("values", "A comma separated list of constants and their value, e.g. a=1,b=2,c=3.").setDefaultValueString("").build()).build()); this->addOption(storm::settings::OptionBuilder(moduleName, statisticsOptionName, false, "Sets whether to display statistics if available.").setShortName(statisticsOptionShortName).build()); + this->addOption(storm::settings::OptionBuilder(moduleName, cudaOptionName, false, "Sets whether to use CUDA to speed up computation time.").build()); } bool GeneralSettings::isHelpSet() const { @@ -288,6 +290,10 @@ namespace storm { bool GeneralSettings::isBisimulationSet() const { return this->getOption(bisimulationOptionName).getHasOptionBeenSet(); } + + bool GeneralSettings::isCudaSet() const { + return this->getOption(cudaOptionName).getHasOptionBeenSet(); + } } // namespace modules } // namespace settings diff --git a/src/settings/modules/GeneralSettings.h b/src/settings/modules/GeneralSettings.h index c65ab57a2..9281d2b57 100644 --- a/src/settings/modules/GeneralSettings.h +++ b/src/settings/modules/GeneralSettings.h @@ -328,6 +328,13 @@ namespace storm { * @return True iff the option was set. */ bool isBisimulationSet() const; + + /*! + * Retrieves whether the option to use CUDA is set. + * + * @return True iff the option was set. + */ + bool isCudaSet() const; bool check() const override; @@ -372,6 +379,7 @@ namespace storm { static const std::string statisticsOptionShortName; static const std::string bisimulationOptionName; static const std::string bisimulationOptionShortName; + static const std::string cudaOptionName; }; } // namespace modules