Merge pull request ComputationalRadiationPhysics#2522 from sbastrakov/fix-PMaccConstexprCaptureMSVC

Add a workaround for MSVC bug with capturing constexpr

Author: psychocoderHPC

include/pmacc/fields/operations/AddExchangeToBorder.hpp

@@ -85,7 +85,7 @@ namespace operations
             // number of cells in a superCell
             constexpr uint32_t numCells = pmacc::math::CT::volume< SuperCellSize >::type::value;
             constexpr uint32_t numWorkers = T_numWorkers;
-            constexpr int dim = T_Mapping::Dim;
+            PMACC_CONSTEXPR_CAPTURE int dim = T_Mapping::Dim;
 
             uint32_t const workerIdx = threadIdx.x;
 

include/pmacc/fields/operations/CopyGuardToExchange.hpp

@@ -82,7 +82,7 @@ namespace operations
             // number of cells in a superCell
             constexpr uint32_t numCells = pmacc::math::CT::volume< SuperCellSize >::type::value;
             constexpr uint32_t numWorkers = T_numWorkers;
-            constexpr int dim = T_Mapping::Dim;
+            PMACC_CONSTEXPR_CAPTURE int dim = T_Mapping::Dim;
 
             uint32_t const workerIdx = threadIdx.x;
 

include/pmacc/particles/ParticlesBase.kernel

@@ -135,7 +135,7 @@ struct KernelFillGapsLastFrame
                 uint32_t const
             )
             {
-                lastFrame = pb.getLastFrame( DataSpace< dim >( superCellIdx ) );
+                lastFrame = pb.getLastFrame( superCellIdx );
                 counterGaps = 0;
                 counterParticles = 0;
                 srcGap = 0;
@@ -324,8 +324,8 @@ struct KernelFillGaps
                 uint32_t const
             )
             {
-                firstFrame = pb.getFirstFrame( DataSpace< dim >(superCellIdx) );
-                lastFrame = pb.getLastFrame( DataSpace< dim >(superCellIdx) );
+                firstFrame = pb.getFirstFrame( superCellIdx );
+                lastFrame = pb.getLastFrame( superCellIdx );
             }
         );
 
@@ -512,7 +512,7 @@ struct KernelShiftParticles
         using FrameType = typename ParBox::FrameType;
         using FramePtr = typename ParBox::FramePtr;
 
-        constexpr uint32_t dim = Mapping::Dim;
+        PMACC_CONSTEXPR_CAPTURE uint32_t dim = Mapping::Dim;
         constexpr uint32_t frameSize = math::CT::volume< typename FrameType::SuperCellSize >::type::value;
         /* number exchanges in 2D=9 and in 3D=27 */
         constexpr uint32_t numExchanges = traits::NumberOfExchanges< dim >::value;
@@ -962,7 +962,7 @@ struct KernelCopyGuardToExchange
         using namespace particles::operations;
         using namespace mappings::threads;
 
-        constexpr uint32_t dim = T_Mapping::Dim;
+        PMACC_CONSTEXPR_CAPTURE uint32_t dim = T_Mapping::Dim;
         constexpr uint32_t frameSize = math::CT::volume< typename T_ParBox::FrameType::SuperCellSize >::type::value;
         constexpr uint32_t numWorkers = T_numWorkers;
 
@@ -1180,7 +1180,7 @@ struct KernelInsertParticles
         using namespace particles::operations;
         using namespace mappings::threads;
 
-        constexpr uint32_t dim = T_Mapping::Dim;
+        PMACC_CONSTEXPR_CAPTURE uint32_t dim = T_Mapping::Dim;
         constexpr uint32_t frameSize = math::CT::volume< typename T_ParBox::FrameType::SuperCellSize >::type::value;
         constexpr uint32_t numWorkers = T_numWorkers;
 

include/pmacc/types.hpp

@@ -232,4 +232,28 @@ enum AreaType
 #define __delete(var) if((var)) { delete (var); var=nullptr; }
 #define __deleteArray(var) if((var)) { delete[] (var); var=nullptr; }
 
+/**
+ * Visual Studio has a bug with constexpr variables being captured in lambdas as
+ * non-constexpr variables, causing build errors. The issue has been verified
+ * for versions 14.0 and 15.5 (latest at the moment) and is also reported in
+ * https://stackoverflow.com/questions/28763375/using-lambda-captured-constexpr-value-as-an-array-dimension
+ * and related issue
+ * https://developercommunity.visualstudio.com/content/problem/1997/constexpr-not-implicitely-captured-in-lambdas.html
+ *
+ * As a workaround (until this is fixed in VS) add a new PMACC_CONSTEXPR_CAPTURE
+ * macro for declaring constexpr variables that are captured in lambdas and have
+ * to remain constexpr inside a lambda e.g., used as a template argument. Such
+ * variables have to be declared with PMACC_CONSTEXPR_CAPTURE instead of
+ * constexpr. The macro will be replaced with just constexpr for other compilers
+ * and for Visual Studio with static constexpr, which makes it capture properly.
+ *
+ * Note that this macro is to be used only in very few cases, where not only a
+ * constexpr is captured, but also it has to remain constexpr inside a lambda.
+ */
+#ifdef _MSC_VER
+#   define PMACC_CONSTEXPR_CAPTURE static constexpr
+#else
+#   define PMACC_CONSTEXPR_CAPTURE constexpr
+#endif
+
 } //namespace pmacc