From 5aa0445f8fc593ccaf3455e318fbaa30a0581fbe Mon Sep 17 00:00:00 2001
From: Martin Bauer <martin.bauer@fau.de>
Date: Wed, 24 Apr 2019 13:03:26 +0200
Subject: [PATCH] Improvements for GPU code generation
- turned on restrict keyword by default (makes large difference on GPUs)
- smarter block indexing: changing block size depending on domain size
Example: previously there where (1,1,1) blocks when requested
block size was (64, 1, 1) and domain size (1, 512, 512), now the
block size is changed automatically to (1, 64, 1) in this case
- added __lauch_bounds__ to kernels to allow better optimizations from
the CUDA compiler
---
.../templates/LatticeModel.tmpl.cpp | 29 +++++++------------
1 file changed, 10 insertions(+), 19 deletions(-)
diff --git a/lbmpy_walberla/templates/LatticeModel.tmpl.cpp b/lbmpy_walberla/templates/LatticeModel.tmpl.cpp
index 5671bde..6f80084 100644
--- a/lbmpy_walberla/templates/LatticeModel.tmpl.cpp
+++ b/lbmpy_walberla/templates/LatticeModel.tmpl.cpp
@@ -39,18 +39,13 @@
#define FUNC_PREFIX __global__
{%- endif %}
-#ifdef WALBERLA_CXX_COMPILER_IS_GNU
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Wunused-variable"
-#pragma GCC diagnostic ignored "-Wunused-parameter"
-#pragma GCC diagnostic ignored "-Wshadow"
-#endif
-
-#ifdef WALBERLA_CXX_COMPILER_IS_CLANG
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wunused-variable"
-#pragma clang diagnostic ignored "-Wunused-parameter"
-#pragma clang diagnostic ignored "-Wshadow"
+#if ( defined WALBERLA_CXX_COMPILER_IS_GNU ) || ( defined WALBERLA_CXX_COMPILER_IS_CLANG )
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Wfloat-equal"
+# pragma GCC diagnostic ignored "-Wshadow"
+# pragma GCC diagnostic ignored "-Wconversion"
+# pragma GCC diagnostic ignored "-Wunused-variable"
+# pragma GCC diagnostic ignored "-Wunused-parameter"
#endif
@@ -128,10 +123,6 @@ mpi::RecvBuffer & operator>> (mpi::RecvBuffer & buf, ::walberla::{{namespace}}::
} // namespace mpi
} // namespace walberla
-#ifdef WALBERLA_CXX_COMPILER_IS_GNU
-#pragma GCC diagnostic pop
-#endif
-
-#ifdef WALBERLA_CXX_COMPILER_IS_CLANG
-#pragma clang diagnostic pop
-#endif
+#if ( defined WALBERLA_CXX_COMPILER_IS_GNU ) || ( defined WALBERLA_CXX_COMPILER_IS_CLANG )
+# pragma GCC diagnostic pop
+#endif
\ No newline at end of file
--
GitLab