Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Using KOKKOS_CLASS_LAMBDA in a class with Kokkos::Random_XorShift64_Pool member data #1696

Closed
Char-Aznable opened this issue Jul 11, 2018 · 5 comments
Assignees
Labels
Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Milestone

Comments

@Char-Aznable
Copy link
Contributor

Char-Aznable commented Jul 11, 2018

It seems to me that in order to use KOKKOS_CLASS_LAMBDA on device code, say CUDA, the class copy constructor has to be KOKKOS_FUNCTION. But the member Kokkos::Random_XorShift64_Pool data doesn't have KOKKOS_FUNCTION copy ctor so I'm stuck. Consider the following code for example:

#include <Kokkos_Core.hpp>
#include <Kokkos_DualView.hpp>
#include <Kokkos_Random.hpp>
#include <iostream>

using namespace std;

using ExecutionSpace = Kokkos::DefaultExecutionSpace;

struct Foo {

  KOKKOS_FUNCTION
  Foo(uint64_t rndSeed) : rndPool_(rndSeed) {};

  template < class E >
  KOKKOS_INLINE_FUNCTION void move() {
    
    using TeamPolicy = Kokkos::TeamPolicy<E>;
    using Team = typename TeamPolicy::member_type;
    TeamPolicy policy(100, 256);


    Kokkos::parallel_for(policy, KOKKOS_CLASS_LAMBDA (Team team) {
        auto rndEngine = rndPool_.get_state();
        int i = rndEngine.urand(0,10);
      }); 
  }

  Kokkos::Random_XorShift64_Pool<ExecutionSpace> rndPool_;
};



int main(int argc, char* argv[]) {
  Kokkos::initialize(argc,argv); {

  Foo d(829819);

  d.move<ExecutionSpace>();

  Kokkos::fence();

  } Kokkos::finalize();
}

This gives error in clang 6 + cuda 9:

testlambda.cpp:23:34: error: reference to __host__ function 'Foo' in __host__ __device__ function
    Kokkos::parallel_for(policy, KOKKOS_CLASS_LAMBDA (Team team) {

Is there a way around this? @ibaned I saw your https://github.com/ibaned/lambda_users_guide post but am I missing something here?

@Char-Aznable
Copy link
Contributor Author

I've tried adding a copy ctor KOKKOS_FUNCTION Foo(const Foo& src) : rndPool_(src.rndPool_) {}; but this triggers the error:

testlambda.cpp:74:41: error: reference to __host__ function 'Random_XorShift64_Pool' in __host__ __device__ function
  KOKKOS_FUNCTION Foo(const Foo& src) : rndPool_(src.rndPool_) {};
                                        ^

@crtrott
Copy link
Member

crtrott commented Jul 12, 2018

Ok this has something to do with the way how clang handles implicitly defaulted constructors. Basically you need to add the default constructors to your class:

  KOKKOS_FUNCTION
  Foo() = default;
  
  KOKKOS_FUNCTION
  Foo(const Foo&) = default;

And we need to add KOKKOS_INLINE_FUNCTION to the random pool classes for the default, the copy constructor and the assignment operator, but not the constructor which takes a seed.

@crtrott crtrott self-assigned this Jul 12, 2018
@crtrott crtrott added the Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos) label Jul 12, 2018
@crtrott crtrott added this to the 2018 July milestone Jul 12, 2018
@crtrott
Copy link
Member

crtrott commented Jul 16, 2018

Just so you know I got a fix as part of some other thing I have been working on. Trying to get that in by tomorrow.

@crtrott
Copy link
Member

crtrott commented Aug 13, 2018

Should be fixed in develop

@Char-Aznable
Copy link
Contributor Author

The patch in develop branch works with or without explicitly marking Foo's default ctor KOKKOS_FUNCTION. Thanks a lot! @crtrott

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Projects
None yet
Development

No branches or pull requests

2 participants