Saitama10000
Saitama10000

Reputation: 286

Mixing cuda and cpp templates and lambdas

Example code: https://github.com/Saitama10000/Mixing-cuda-and-cpp-templates-and-lambdas

The example code doesn't compile. I am supposed to add a template instantiation in the .cu file, but I don't know how. I've tried this:

typedef float(*op)(float);
template std::vector<float> f<op>(std::vector<float> const&, op);

but I still get this compilation error:

In file included from Mixing-cuda-and-cpp-templates-and-lambdas/main.cpp:6:
Mixing-cuda-and-cpp-templates-and-lambdas/kernel.cuh:6:20: error: ‘std::vector<float> f(const std::vector<float>&, FUNC) [with FUNC = main()::<lambda(float)>]’, declared using local type ‘main()::<lambda(float)>’, is used but never defined [-fpermissive]
    6 | std::vector<float> f(std::vector<float> const& a, FUNC func);
      |                    ^
Mixing-cuda-and-cpp-templates-and-lambdas/kernel.cuh:6:20: warning: ‘std::vector<float> f(const std::vector<float>&, FUNC) [with FUNC = main()::<lambda(float)>]’ used but never defined
make[2]: *** [CMakeFiles/main.dir/build.make:82: CMakeFiles/main.dir/main.cpp.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:95: CMakeFiles/main.dir/all] Error 2
make: *** [Makefile:103: all] Error 2

Upvotes: 0

Views: 169

Answers (1)

Abator Abetor
Abator Abetor

Reputation: 2598

There are two problems with your approach.

First, each lambda has it's own type even if parameters and function body are the same.

For example, the following assertion fails

#include <type_traits>

int main(){
    auto lambda1 = [](){};
    auto lambda2 = [](){};

    static_assert(std::is_same<decltype(lambda1), decltype(lambda2)>::value, "not same");
}

That means, even if you somehow manage to explitely instantiate your template with the type of the lambda it won't be the type of the lambda which you will pass to your function. This problem can be solved by using functors instead of lambdas. Define a set of functors which may be used to call the function, and use them for template instantiation.

Second, you want to pass a __host__ __device__ function. This annotation is a CUDA C++ extension which cannot be compiled with a standard C++ compiler. You have to use a CUDA compiler instead, which in turn allows you to place your kernel and wrappers in the .cuh file.

Upvotes: 2

Related Questions