Error building with Cuda when passing -DKOKKOS_CUDA_USE_LAMBDA to generate_makefile.bash

[ndellingwood]

I am getting compile errors when running make build-test with Cuda (from current develop branch) if I generate a makefile and pass -DKOKKOS_CUDA_USE_LAMBDA to cxxflags. This used to work for me, is it no longer correct to use this? If I don't pass this to cxxflags I have no problems.

The errors reported (pasted at end) occur in TestReduce.hpp with instantiations of ReduceCombinatorical::

Currently Loaded Modulefiles:

  1) git/2.1.3                                  4) gcc/4.9.2/openmpi/1.8.7/cuda/7.5.18
  2) cmake/2.8.11                          5) cuda/7.5.18
  3) gcc/4.9.2/base

Makefile that caused build errors generated like so:
~/kokkos/generate_makefile.bash --with-cuda --compiler=~/kokkos/config/nvcc_wrapper --cxxflags="-expt-extended-lambda -DKOKKOS_CUDA_USE_LAMBDA"

Removing -DKOKKOS_CUDA_USE_LAMBDA works fine:
~/kokkos/generate_makefile.bash --with-cuda --compiler=~/kokkos/config/nvcc_wrapper --cxxflags="-expt-extended-lambda"

Sample error (repeats multiple times)

/ascldap/users/ndellin/kokkos/core/src/Cuda/Kokkos_CudaExec.hpp(267): error: A type defined inside a __host__ function  ("lambda [](const int &, double &)->void") cannot be used in the template argument type of a __global__ function template instantiation
          detected during:
            instantiation of "Kokkos::Impl::cuda_parallel_launch_local_memory" based on template argument <Kokkos::Impl::ParallelReduce<Kokkos::Impl::CudaFunctorAdapter<lambda [](const int &, double &)->void, Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, double, void>, Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, Kokkos::InvalidType, Kokkos::CudaSpace::execution_space>>
(267): here
            instantiation of "Kokkos::Impl::CudaParallelLaunch<DriverType, false>::CudaParallelLaunch(const DriverType &, const dim3 &, const dim3 &, int, cudaStream_t) [with DriverType=Kokkos::Impl::ParallelReduce<Kokkos::Impl::CudaFunctorAdapter<lambda [](const int &, double &)->void, Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, double, void>, Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, Kokkos::InvalidType, Kokkos::CudaSpace::execution_space>]"
/ascldap/users/ndellin/kokkos/core/src/Cuda/Kokkos_Cuda_Parallel.hpp(829): here
            instantiation of "void Kokkos::Impl::ParallelReduce<FunctorType, Kokkos::RangePolicy<Traits...>, ReducerType, Kokkos::Cuda>::execute() [with FunctorType=Kokkos::Impl::CudaFunctorAdapter<lambda [](const int &, double &)->void, Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, double, void>, ReducerType=Kokkos::InvalidType, Traits=<Kokkos::CudaSpace::execution_space>]"
/ascldap/users/ndellin/kokkos/core/src/Kokkos_Parallel_Reduce.hpp(250): here
            instantiation of "void Kokkos::Impl::ParallelReduceAdaptor<PolicyType, FunctorType, ReturnType>::execute(const std::string &, const PolicyType &, const FunctorType &, ReturnType &) [with PolicyType=Kokkos::RangePolicy<Kokkos::CudaSpace::execution_space>, FunctorType=lambda [](const int &, double &)->void, ReturnType=double]"
/ascldap/users/ndellin/kokkos/core/src/Kokkos_Parallel_Reduce.hpp(339): here
            instantiation of "void Kokkos::parallel_reduce(const size_t &, const FunctorType &, ReturnType &) [with FunctorType=lambda [](const int &, double &)->void, ReturnType=double]"
/ascldap/users/ndellin/kokkos/core/unit_test/TestReduce.hpp(838): here
            instantiation of "void Test::ReduceCombinatorical::AddReturnArgument(Args...) [with Args=<int, lambda [](const int &, double &)->void>]"
/ascldap/users/ndellin/kokkos/core/unit_test/TestReduce.hpp(875): here
            instantiation of "void Test::ReduceCombinatorical::AddLambdaRange(Args...) [with Args=<int>]" 
/ascldap/users/ndellin/kokkos/core/unit_test/TestReduce.hpp(921): here
            instantiation of "void Test::ReduceCombinatorical::AddFunctorLambdaRange(Args...) [with Args=<int>]"
/ascldap/users/ndellin/kokkos/core/unit_test/TestReduce.hpp(938): here
            instantiation of "void Test::ReduceCombinatorical::AddPolicy(Args...) [with Args=<>]"
/ascldap/users/ndellin/kokkos/core/unit_test/TestReduce.hpp(955): here

[ndellingwood]

@crtrott
In TestReduce.hpp I replaced #ifdef KOKKOS_HAVE_CXX11_DISPATCH_LAMBDA
with
#if defined(KOKKOS_HAVE_CXX11_DISPATCH_LAMBDA) && !defined(KOKKOS_HAVE_CUDA)
in the two functions
void AddFunctorLambdaRange(Args... args) {
and
void AddFunctorLambdaTeam(Args... args) {

Things are compiling again and passing unit tests when using a makefile generated with a cxxflags argument -DKOKKOS_CUDA_USE_LAMBDA, if this is a sufficient fix let me know and I can run more extensive testing and issue a pull request.

[crtrott]

No we need to properly fix this. This looks as if KOKKOS_LAMBDA was not correctly defined as device

[ndellingwood]

Regarding makefile generation, here is the proper way to generate and enable lambdas from our discussion:
~/kokkos/generate_makefile.bash --with-cuda --compiler=~/kokkos/config/nvcc_wrapper --with-cuda-options=enable_lambda

[crtrott]

closing the issue, since it was a configuration error.