Add CUDA LTO.

CMakeLists.txt

@@ -94,12 +94,21 @@ if(BUILD_VF_GPU)
 
     set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
 
-    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcudafe --display_error_number")
 
-    if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-        message(WARNING "CMAKE_CUDA_ARCHITECTURES was not defined and is set to 30 (CUDA support until 10.1 only).")
-        set(CMAKE_CUDA_ARCHITECTURES 30)
-    endif()
+    #set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -dlto")
+    enable_language(CUDA)
+    #set(CMAKE_CUDA_FLAGS "-gencode arch=compute_70,code=lto_70")
+
+    set(CMAKE_CUDA_FLAGS "-dlto -arch=sm_75")
+    #set(CMAKE_CUDA_ARCHITECTURES 75)
+    #set(CMAKE_CUDA_FLAGS "-arch=sm_75")
+
+    #set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcudafe --display_error_number")
+
+    #if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+    #    message(WARNING "CMAKE_CUDA_ARCHITECTURES was not defined and is set to 30 (CUDA support until 10.1 only).")
+    #    set(CMAKE_CUDA_ARCHITECTURES 30)
+    #endif()
 
     message("CUDA Architecture: ${CMAKE_CUDA_ARCHITECTURES}")
 endif()

apps/gpu/LBM/DrivenCavity/CMakeLists.txt

@@ -6,4 +6,8 @@ vf_add_library(BUILDTYPE binary PRIVATE_LINK basics VirtualFluids_GPU GridGenera
 
 set_source_files_properties(DrivenCavity.cpp PROPERTIES LANGUAGE CUDA)
 
-set_target_properties(DrivenCavity PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(DrivenCavity PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
\ No newline at end of file
+
+#target_compile_options(DrivenCavity PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-dlto>)
+
+#target_link_options(DrivenCavity PRIVATE $<DEVICE_LINK:-dlto>)
\ No newline at end of file

src/gpu/GridGenerator/CMakeLists.txt

@@ -9,7 +9,7 @@ set_target_properties(${library_name} PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
 # according to linker error when building static libraries.
 # https://stackoverflow.com/questions/50033435/cmake-cuda-separate-compilation-static-lib-link-error-on-windows-but-not-on-ubun
 if (NOT BUILD_SHARED_LIBRARY)
-    set_target_properties(${library_name} PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON)
+    #set_target_properties(${library_name} PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 endif()
 
 # we want to suppress all cuda warnings so far for this library.

src/gpu/VirtualFluids_GPU/CMakeLists.txt

@@ -13,3 +13,9 @@ linkBoost(COMPONENTS "serialization")
 #SET(TPN_WIN32 "/EHsc")
 #https://stackoverflow.com/questions/6832666/lnk2019-when-including-asio-headers-solution-generated-with-cmake
 #https://stackoverflow.com/questions/27442885/syntax-error-with-stdnumeric-limitsmax
+
+set_target_properties(VirtualFluids_GPU PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+#target_compile_options(VirtualFluids_GPU PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-dlto>)
+
+#target_link_options(VirtualFluids_GPU PRIVATE $<DEVICE_LINK:-dlto>)

src/lbm/CMakeLists.txt

-project(lbm LANGUAGES CXX)
+project(lbm LANGUAGES CXX CUDA)
 
 vf_add_library(NAME lbm PUBLIC_LINK basics)
 
 vf_add_tests()
+
+set_target_properties(lbm PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

src/lbm/CumulantChimeraPreCompiled.h

+#ifndef LBM_CUMULANT_CHIMERA_PRE_H
+#define LBM_CUMULANT_CHIMERA_PRE_H
+
+#ifndef __host__
+#define __host__
+#endif
+#ifndef __device__
+#define __device__
+#endif
+
+#include <cmath>
+
+#include <basics/Core/DataTypes.h>
+#include <basics/Core/RealConstants.h>
+
+#include "D3Q27.h"
+#include "Chimera.h"
+#include "MacroscopicQuantities.h"
+
+namespace VF
+{
+namespace LBM
+{
+
+struct Distribution27
+{
+    real f[27];
+
+    inline __host__ __device__ real getDensity_() const
+    {
+        return getDensity(f);
+    }
+};
+
+
+
+
+inline __host__ __device__ real abs_internal(real value)
+{
+#ifdef __CUDA_ARCH__
+    return ::abs(value);
+#else
+    return std::abs(value);
+#endif
+}
+
+
+
+//////////////////////////////////////////////////////////////////////////
+//! Cumulant K17 Kernel is based on \ref
+//! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 ]</b></a>
+//! and \ref
+//! <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.07.004 ]</b></a>
+//////////////////////////////////////////////////////////////////////////
+__host__ __device__ void cumulantChimera(Distribution27& distribution, real omega, real* forces);
+
+}
+}
+#endif