Tell Trilinos packages about host+device lambda

[ibaned]

Please see the discussion in trilinos/Trilinos#1278 about how to manage the move away from CUDA 7.5 with #ifdefs.

This issue in particular is my proposal that Kokkos define some preprocessor symbol in the cases where a KOKKOS_LAMBDA will work in all enabled execution spaces, as determined by the existing logic in Kokkos_Macros.hpp. Downstream Trilinos packages can then #ifdef their code based on this symbol, and handle the negative case with functors during the remaining period of support for CUDA 7.5, while being able to move forward with new CUDA 8.0-style lambda-based development.

[hcedwar]

We have a naming & definition inconsistency with our LAMBDA macros that needs to be fixed.
Issue #150 defined some guidelines.

[hcedwar]

The meaning of the KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA output macro should be "dispatch lambda is available for all configured back-ends / execution spaces."
Currently means "is available for the default execution space."

[mhoemmen]

@hcedwar @ibaned Does this mean that if I'm building with CUDA 7.5, that I now need to set a nondefault CMake option in order for KOKKOS_LAMBDA to include the __device__ marking? Thanks!

[ibaned]

Nope, it doesn't change the behavior of KOKKOS_LAMBDA, which for the record is as follows:

1. Outside CUDA, just [=]
2. With CUDA 7.5, [=]__device__ as long as Kokkos_ENABLE_Cuda_Lambda=ON in CMake
3. With CUDA 8.0, [=]__host__ __device__ as long as Kokkos_ENABLE_Cuda_Lambda=ON in CMake

[mhoemmen]

@ibaned Just to clarify: Before, with CUDA 7.5, KOKKOS_LAMBDA was [=]__device__, even if I did not explicitly set Kokkos_ENABLE_Cuda_Lambda=ON in CMake. Is that still the case, or do I now need to set that variable explicitly? Sorry to be so specific :-) .

[ibaned]

Specific is good :)
Unfortunately, I don't think this is true, it is OFF by default (based on my inspection of the code, see below)
https://github.com/kokkos/kokkos/blob/master/CMakeLists.txt#L75
https://github.com/kokkos/kokkos/blob/master/core/cmake/KokkosCore_config.h.in#L60
https://github.com/kokkos/kokkos/blob/master/core/src/Kokkos_Macros.hpp#L126
We could open another issue to enable this feature by default, which I would fully support.

Edit: I haven't actually tested this with a real configuration

[ibaned]

So PR #1020 tried to fix this but introduced a warning in Trilinos builds:

In file included from /home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Core_fwd.hpp:51:0,
                 from /home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Core.hpp:50,
                 from /home/daibane/src/Trilinos/packages/teuchos/kokkoscompat/src/KokkosCompat_View.hpp:58,
                 from /home/daibane/src/Trilinos/packages/teuchos/kokkoscompat/src/KokkosCompat_ClassicNodeAPI_Wrapper.hpp:5,
                 from /home/daibane/src/Trilinos/packages/tpetra/classic/NodeAPI/Kokkos_DefaultNode.hpp:48,
                 from /home/daibane/src/Trilinos/packages/tpetra/core/src/Tpetra_ConfigDefs.hpp:46,
                 from /home/daibane/src/Trilinos/packages/tpetra/core/src/Tpetra_Map_decl.hpp:49,
                 from /home/daibane/build/gcc/Trilinos/packages/tpetra/core/src/Tpetra_Map.hpp:1,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraVectorSpace_decl.hpp:48,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraLinearOp_decl.hpp:46,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraLinearOp.hpp:44,
                 from /home/daibane/src/Trilinos/packages/ifpack2/adapters/thyra/Thyra_Ifpack2PreconditionerFactory_def.hpp:48,
                 from /home/daibane/build/gcc/Trilinos/packages/ifpack2/adapters/Thyra_Ifpack2PreconditionerFactory.hpp:2,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_InvertMassMatrixDecorator_Def.hpp:57,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_InvertMassMatrixDecorator.hpp:165,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_RythmosSolver_Def.hpp:76,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_RythmosSolver.hpp:203,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_Epetra_RythmosSolver.hpp:53,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_Epetra_RythmosSolver.cpp:43:
/home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Macros.hpp:105:0: warning: "KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA" redefined
 #define KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA
 
In file included from /home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Macros.hpp:65:0,
                 from /home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Core_fwd.hpp:51,
                 from /home/daibane/src/Trilinos/kokkos/core/src/Kokkos_Core.hpp:50,
                 from /home/daibane/src/Trilinos/packages/teuchos/kokkoscompat/src/KokkosCompat_View.hpp:58,
                 from /home/daibane/src/Trilinos/packages/teuchos/kokkoscompat/src/KokkosCompat_ClassicNodeAPI_Wrapper.hpp:5,
                 from /home/daibane/src/Trilinos/packages/tpetra/classic/NodeAPI/Kokkos_DefaultNode.hpp:48,
                 from /home/daibane/src/Trilinos/packages/tpetra/core/src/Tpetra_ConfigDefs.hpp:46,
                 from /home/daibane/src/Trilinos/packages/tpetra/core/src/Tpetra_Map_decl.hpp:49,
                 from /home/daibane/build/gcc/Trilinos/packages/tpetra/core/src/Tpetra_Map.hpp:1,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraVectorSpace_decl.hpp:48,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraLinearOp_decl.hpp:46,
                 from /home/daibane/src/Trilinos/packages/thyra/adapters/tpetra/src/Thyra_TpetraLinearOp.hpp:44,
                 from /home/daibane/src/Trilinos/packages/ifpack2/adapters/thyra/Thyra_Ifpack2PreconditionerFactory_def.hpp:48,
                 from /home/daibane/build/gcc/Trilinos/packages/ifpack2/adapters/Thyra_Ifpack2PreconditionerFactory.hpp:2,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_InvertMassMatrixDecorator_Def.hpp:57,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_InvertMassMatrixDecorator.hpp:165,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_RythmosSolver_Def.hpp:76,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_RythmosSolver.hpp:203,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_Epetra_RythmosSolver.hpp:53,
                 from /home/daibane/src/Trilinos/packages/piro/src/Piro_Epetra_RythmosSolver.cpp:43:
/home/daibane/src/Trilinos/kokkos/core/src/impl/Kokkos_OldMacros.hpp:145:0: note: this is the location of the previous definition
 #define KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA KOKKOS_HAVE_CXX11_DISPATCH_LAMBDA

I guess this flag used to have meaning as input, back when C++11 was optional. This warning blocks promotion, even though its not a compile error or runtime issue, we'll look very bad if we let it get into Trilinos (also it makes the meaning of the output flag wrong). Will open a PR momentarily to remove any input meaning of KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA, since C++11 has been required for a while.