kokkos / llvm-project

The LLVM Project is a collection of modular and reusable compiler and toolchain technologies. Note: the repository does not accept github pull requests at this moment. Please submit your patches at http://reviews.llvm.org.
http://llvm.org
Other
4 stars 2 forks source link

Kokkos ensure kokkos function #16

Open calewis opened 4 years ago

calewis commented 4 years ago

The ensure-kokkos-function check will catch the following issues in users code:

void fooOOPS(int i){printf("%i",i);}

KOKKOS_FUNCTION void bar(int i){
   fooOOPS(i); // Should warn
}

void bar2(){
  Kokkos::parallel_for(15, KOKKOS_LAMBDA(int i){
      fooOOPS(i); // Should warn
      });
}

In both of these examples it is possible that the code will compile and run acceptably in certain circumstances, default execution space is OpenMP or other host, but will fail when compiled for the cuda backend. The goal is to help users who mainly develop on host but then need to run on devices.

We have provided a way to opt out of the check when the default host execution space is used.

void fooOOPS(int i){printf("%i",i);}

// We require the annotation to turn off the warning below, 
// this annotation is added to Kokkos::DefaultHostExecutionSpace automatically in develop
using MyHostSpace
#if defined(__clang_analyzer__)
 [[clang::annotate("DefaultHostExectutionSpace")]] 
#endif
= Kokkos::Serial;

void bar2(){
  Kokkos::parallel_for(Kokkos::RangePolicy<MyHostSpace>(15), KOKKOS_LAMBDA(int i){
      fooOOPS(i); // This warning can be turned off, but is on by default.  
      });
}

There are a few minor corner cases that might trigger the warning or will be false negatives, but I deemed the work to catch them to be excessive without complaints from users

template<typename F>
KOKKOS_FUNCTION void runF(F f){
  f(); // Won't warn if this is a lambda even if it breaks on cuda
}

void bar(){
  auto func = []{};
  Kokkos::parallel_for(15, KOKKOS_LAMBDA(int i){
      runF(func); // Won't warn
    });
}

The reason this won't warn is because we have to allow the following

KOKKOS_FUNCTION void foo(){
  []{}(); // Can't warn here, detecting that lambda::operator() is local to foo is not trivial. 
}

I'm sure that with enough users there will be other issues, but I think it's robust enough for release.

calewis commented 4 years ago

Closes #10

dalg24 commented 4 years ago

If a user relies on Clang's implementation and overloads a function on __host__ and __device__ this will get flagged right? (Not that I think it is a big deal, just checking that I understand)

calewis commented 4 years ago

If a user relies on Clang's implementation and overloads a function on __host__ and __device__ this will get flagged right? (Not that I think it is a big deal, just checking that I understand)

Yes, if a user writes functions that are legal in the context they call them and they are careful to not call them in other contexts then they will get warnings. Is that likely (this is the reason for letting people opt out if they explicitly use the default host space)? Also I suspect that these users might not be interested in a check that requires them to use KOKKOS_FUNCTION

dalg24 commented 3 years ago

Any reason to delay the merge?