add a cuda test and make sure it actually works

Author: hschwane

CMakeLists.txt

@@ -12,7 +12,7 @@ include(GNUInstallDirs)
 include(GetVersionFromGit)
 
 # possibility to turn of features with special dependencies
-option(USE_CUDA "Compile with cuda features" OFF)
+option(USE_CUDA "Compile with cuda features" ON)
 option(USE_OPENGL "Compile with openGL features" ON)
 
 # enable languages
@@ -115,6 +115,9 @@ target_link_libraries( mpUtils  ${LIBRARIES})
 
 # set properties
 set_target_properties(mpUtils PROPERTIES VERSION ${VERSION_SHORT} SOVERSION ${VERSION_MAJOR})
+if (USE_CUDA)
+    set_target_properties( mpUtils PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+endif ()
 
 # set include dir (in this case also the src)
 target_include_directories(mpUtils PUBLIC src ${INCLUDE_PATHES})

exec/devTest/CMakeLists.txt

@@ -1,12 +1,13 @@
-cmake_minimum_required(VERSION 3.3)
+cmake_minimum_required(VERSION 3.8)
 
 # set src files
 set(SOURCE_FILES
-    main.cpp
-        )
+        main.cpp
+        cudaTest.cu)
 
 # create target
 add_executable(devTest ${SOURCE_FILES})
+#set_target_properties( devTest PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
 
 # link libraries
 target_link_libraries(devTest mpUtils)

exec/devTest/cudaTest.cu

@@ -0,0 +1,88 @@
+/*
+ * mpUtils
+ * cudaTest.cpp
+ *
+ * @author: Hendrik Schwanekamp
+ * @mail:   [email protected]
+ *
+ * Copyright (c) 2018 Hendrik Schwanekamp
+ *
+ */
+
+#include "cudaTest.h"
+#include <Cuda/cudaUtils.h>
+
+class Managed
+{
+public:
+    void *operator new(size_t len) {
+        void *ptr;
+        assert_cuda(cudaMallocManaged(&ptr, len));
+        assert_cuda(cudaDeviceSynchronize());
+        return ptr;
+    }
+
+    void operator delete(void *ptr) {
+        assert_cuda(cudaDeviceSynchronize());
+        assert_cuda(cudaFree(ptr));
+    }
+};
+
+class Array : public Managed
+{
+public:
+    Array(size_t n) : length(n) {cudaMallocManaged(&data,n* sizeof(float));}
+    ~Array() {cudaFree(data);}
+
+    __host__ __device__
+    float& operator[](int pos) {return data[pos];}
+
+    const size_t length;
+private:
+    float* data;
+};
+
+// CUDA kernel to add elements of two arrays
+__global__
+void add(size_t n, Array& x, Array& y)
+{
+    int index = blockIdx.x * blockDim.x + threadIdx.x;
+    int stride = blockDim.x * gridDim.x;
+
+
+    for (int i = index; i < n; i += stride)
+    {
+        y[i] = x[i] + y[i];
+    }
+}
+
+void testCuda()
+{
+    size_t N = 1<<20;
+    auto x=std::make_unique<Array>(N);
+    auto y=std::make_unique<Array>(N);
+
+    assert_cuda( cudaPeekAtLastError() );
+
+    // initialize x and y arrays on the host
+    for (int i = 0; i < N; i++) {
+        (*x)[i] = 1.0f;
+        (*y)[i] = 2.0f;
+    }
+
+    // Launch kernel on 1M elements on the GPU
+    int blockSize = 256;
+    int numBlocks = (N + blockSize - 1) / blockSize;
+    add<<<numBlocks, blockSize>>>(N, *x, *y);
+
+    // Wait for GPU to finish before accessing on host
+    cudaDeviceSynchronize();
+
+    // Check for errors (all values should be 3.0f)
+    float maxError = 0.0f;
+    for (int i = 0; i < N; i++)
+        maxError = fmax(maxError, fabs((*y)[i]-3.0f));
+    logINFO("TEST") << "Max error: " << maxError << std::endl;
+
+    return;
+}

exec/devTest/cudaTest.h

@@ -0,0 +1,20 @@
+/*
+ * mpUtils
+ * cudaTest.h
+ *
+ * @author: Hendrik Schwanekamp
+ * @mail:   [email protected]
+ *
+ * Copyright (c) 2018 Hendrik Schwanekamp
+ *
+ */
+#ifndef MPUTILS_CUDATEST_H
+#define MPUTILS_CUDATEST_H
+
+// includes
+//--------------------
+//--------------------
+
+void testCuda();
+
+#endif //MPUTILS_CUDATEST_H

exec/devTest/main.cpp

@@ -12,34 +12,14 @@
  *
  */
 
-#include <stringUtils.h>
-#include <Timer/Stopwatch.h>
-#include <iostream>
-#include <Log/Log.h>
-#include <Log/ConsoleSink.h>
-#include <Log/FileSink.h>
-#include <Log/SyslogSink.h>
+#include <mpUtils.h>
 #include <chrono>
-
-#include <typeinfo>
-#include <ctime>
-#include "Timer/DeltaTimer.h"
-#include <thread>
-#include "Timer/Timer.h"
-#include "Timer/AsyncTimer.h"
-#include <thread>             // std::thread
-#include <mutex>              // std::mutex, std::unique_lock
-#include <condition_variable> // std::condition_variable
-#include <syslog.h>
-
+#include "cudaTest.h"
 
 using namespace mpu;
 using namespace std;
 using namespace std::chrono;
 
-constexpr int numRuns = 10;
-double dTime =0;
-
 int main()
 {
     CpuStopwatch timer;
@@ -50,19 +30,10 @@ int main()
 
         myLog(LogLvl::INFO, MPU_FILEPOS , "TEST") << "Hi, a log";
 
-        logINFO("TEST") << "Some generic Info";
-        logWARNING("TEST") << "Some log warning";
-        logERROR("MODULE_TEST") << "some stuff has happend";
-        logDEBUG("some stuff") << "some stuff is debugging stuff";
-        logDEBUG2("some stuff") << "more debugging stuff";
-
-    timer.getSeconds();
-
-    yield();
-    sleep(2);
-    yield();
+        logINFO("TEST") << "Testing Cuda.";
+        testCuda();
 
+    logINFO("TEST") << "It took me " << timer.getSeconds() << " seconds" << endl;
     myLog.close();
-    cout << "It took me " << dTime << " seconds" << endl;
     return 0;
 }

src/Cuda/cudaUtils.h

@@ -19,6 +19,7 @@
 #include <device_launch_parameters.h>
 #include <cstdio>
 #include <cstdlib>
+#include "../Log/Log.h"
 //--------------------
 
 // make clion understand some cuda specific stuff
@@ -55,4 +56,19 @@ typedef long long longlong;
 #endif
 //--------------------
 
+#define assert_cuda(CODE) _cudaAssert((CODE),MPU_FILEPOS);
+
+inline void _cudaAssert(cudaError_t code, std::string&& filepos)
+{
+    if (code != cudaSuccess)
+    {
+        std::string message("Cuda error: " + std::string(cudaGetErrorString(code)));
+
+        if (!(mpu::Log::noGlobal() || mpu::Log::getGlobal().getLogLevel() < mpu::LogLvl::FATAL_ERROR))
+            mpu::Log::getGlobal()(mpu::LogLvl::FATAL_ERROR, std::move(filepos), "cuda") << message;
+
+        throw std::runtime_error("Cuda error: " + message);
+    }
+}
+
 #endif //MPUTILS_CUDAUTILS_H

src/Graphics/Window.h

@@ -90,14 +90,14 @@ class Window
     GLFWwindowiconifyfun setMinimizeCallback(GLFWwindowiconifyfun cb) {glfwSetWindowIconifyCallback(m_w.get(),cb);} //!< callback will be called when the window is minimized
 
     // input callbacks
-    GLFWkeyfun setKeyCallback(GLFWkeyfun cb) {glfwSetKeyCallback(m_w.get(),cb);} //!< callback will be called when key input is availible
-    GLFWcharfun setCharCallback(GLFWcharfun cb) {glfwSetCharCallback(m_w.get(),cb);} //!< callback provides character input
-    GLFWcharmodsfun setCharmodsCallback(GLFWcharmodsfun cb) {glfwSetCharModsCallback(m_w.get(),cb);} //!< callback provides charater input with modifier keys
-    GLFWmousebuttonfun setMousebuttonCallback(GLFWmousebuttonfun cb) {glfwSetMouseButtonCallback(m_w.get(),cb);} //!< called when a mouse button is pressed
-    GLFWcursorposfun setCoursorposCallback(GLFWcursorposfun cb) {glfwSetCursorPosCallback(m_w.get(),cb);} //!< called when the cursor is moved
-    GLFWcursorenterfun setCoursorenterCallback(GLFWcursorenterfun cb) {glfwSetCursorEnterCallback(m_w.get(),cb);} //!< called when the cursor enters or leaves the window
-    GLFWscrollfun setScrollCallback(GLFWscrollfun cb) {glfwSetScrollCallback(m_w.get(),cb);} //!< called when the scroll wheel is moved
-    GLFWdropfun setDropCallbac(GLFWdropfun cb) {glfwSetDropCallback(m_w.get(),cb);} //!< called when someting is drag'n droped onto the window
+    GLFWkeyfun setKeyCallback(GLFWkeyfun cb) {return glfwSetKeyCallback(m_w.get(),cb);} //!< callback will be called when key input is availible
+    GLFWcharfun setCharCallback(GLFWcharfun cb) {return glfwSetCharCallback(m_w.get(),cb);} //!< callback provides character input
+    GLFWcharmodsfun setCharmodsCallback(GLFWcharmodsfun cb) {return glfwSetCharModsCallback(m_w.get(),cb);} //!< callback provides charater input with modifier keys
+    GLFWmousebuttonfun setMousebuttonCallback(GLFWmousebuttonfun cb) {return glfwSetMouseButtonCallback(m_w.get(),cb);} //!< called when a mouse button is pressed
+    GLFWcursorposfun setCoursorposCallback(GLFWcursorposfun cb) {return glfwSetCursorPosCallback(m_w.get(),cb);} //!< called when the cursor is moved
+    GLFWcursorenterfun setCoursorenterCallback(GLFWcursorenterfun cb) {return glfwSetCursorEnterCallback(m_w.get(),cb);} //!< called when the cursor enters or leaves the window
+    GLFWscrollfun setScrollCallback(GLFWscrollfun cb) {return glfwSetScrollCallback(m_w.get(),cb);} //!< called when the scroll wheel is moved
+    GLFWdropfun setDropCallbac(GLFWdropfun cb) {return glfwSetDropCallback(m_w.get(),cb);} //!< called when someting is drag'n droped onto the window
 
     // input functions
     void setInputMode(int mode, int value) {glfwSetInputMode(m_w.get(),mode,value);} //!< see glfwSetInputMode for reference

src/mpUtils.h

@@ -38,10 +38,5 @@
     #include "Graphics/Graphics.h"
 #endif
 
-// include cuda
-#ifdef USE_CUDA
-    #include "Cuda/cudaUtils.h"
-#endif
-
 
 #endif //MPUTILS_MPUTILS_H

src/stringUtils.h

@@ -122,7 +122,7 @@ auto makeFuncCopyable( F&& f )
     auto spf = std::make_shared<F>(std::forward<F>(f) );
     return [spf](auto&&... args)->decltype(auto)
     {
-        return (*spf)( decltype(args)(args)... );
+        return (*spf)( (args)... );
     };
 }