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

Kokkos ensure kokkos function #16

Open
wants to merge 19 commits into
base: master
Choose a base branch
from

Conversation

calewis
Copy link

@calewis calewis commented Jul 15, 2020

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
Copy link
Author

calewis commented Jul 15, 2020

Closes #10

Copy link
Author

@calewis calewis left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Also the checks for implicit this got merged into this PR because I needed the Kokkos mock header and got my wires crossed in getting it.

@@ -83,6 +83,7 @@ buildCompilerInvocation(const ParseInputs &Inputs, clang::DiagnosticConsumer &D,
CI->getPreprocessorOpts().PCHThroughHeader.clear();
CI->getPreprocessorOpts().PCHWithHdrStop = false;
CI->getPreprocessorOpts().PCHWithHdrStopCreate = false;
CI->getPreprocessorOpts().SetUpStaticAnalyzer = true;
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just so we are clear this might be an issue going forward, but without this ensure kokkos function doesn't work in clangd

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please elaborate on why this might become an issue.

Copy link
Author

@calewis calewis Aug 26, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It possibly turns on code paths that wouldn't compile when you are actually running the compiler. If downstream projects use __clang_analyzer__ to do stuff for their projects, this might interfere with their usage of our clangd. The alternative is that we define something like __kokkos_analyzer__ that shouldn't have any conflicts, but I think this is less work until someone complains.

I don't know why clangd doesn't turn this on in the upstream llvm, but I'm going to appeal to https://en.wikipedia.org/wiki/Wikipedia:Chesterton%27s_fence and not investigate farther.

@calewis calewis changed the title [WIP] Kokkos ensure kokkos function Kokkos ensure kokkos function Aug 19, 2020
CA Lewis added 18 commits August 19, 2020 14:53
We might need to make something like a __kokkos_analyzer__ definition in
the future to more carefully target only kokkos code.  I'm open to
suggestions here.
If the user wants to allow certain things that would normally trigger
lint checks when they use
`Kokkos::DefaultHostExectuionSpace` we have to try to detect that
typedef.
I think everything is done except for detecting the methods on user
defined functors. I need to workout which functor methods require
annotation and which ones don't.
KOKKOS_LAMBDA

Also did a bit of clean up so I could reuse code.
This also required the start of a Kokkos Mocking framework which is
going to be a work in progress for awhile I am sure.
If the user wants to allow certain things that would normally trigger
lint checks when they use
`Kokkos::DefaultHostExectuionSpace` we have to try to detect that
typedef.
Lesson learned: I should write tests before checks :P
@calewis calewis force-pushed the kokkos-ensure-kokkos-function branch from 90e99fc to 09eda9e Compare August 19, 2020 21:54
Copy link

@DavidPoliakoff DavidPoliakoff left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The clangd thing is a bit spicy ([insert Hyrum's law argument here]), but I think it's a good idea

dalg24
dalg24 previously requested changes Aug 26, 2020
Copy link
Member

@dalg24 dalg24 left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please provide a description.

clang-tools-extra/docs/ReleaseNotes.rst Outdated Show resolved Hide resolved
@@ -83,6 +83,7 @@ buildCompilerInvocation(const ParseInputs &Inputs, clang::DiagnosticConsumer &D,
CI->getPreprocessorOpts().PCHThroughHeader.clear();
CI->getPreprocessorOpts().PCHWithHdrStop = false;
CI->getPreprocessorOpts().PCHWithHdrStopCreate = false;
CI->getPreprocessorOpts().SetUpStaticAnalyzer = true;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please elaborate on why this might become an issue.

@dalg24
Copy link
Member

dalg24 commented Aug 27, 2020

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)

@dalg24 dalg24 dismissed their stale review August 27, 2020 01:53

A (very nice) description has been added

@calewis
Copy link
Author

calewis commented Aug 27, 2020

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
Copy link
Member

dalg24 commented Mar 3, 2021

Any reason to delay the merge?

crtrott pushed a commit that referenced this pull request Jul 13, 2023
…tput

The crash happens in clang::driver::tools::SplitDebugName when Output is
InputInfo::Nothing. It doesn't happen with standalone clang driver because
output is created in Driver::BuildJobsForActionNoCache.

Example backtrace:
```
* thread #1, name = 'clangd', stop reason = hit program assert
  * frame #0: 0x00007ffff5c4eacf libc.so.6`raise + 271
    frame #1: 0x00007ffff5c21ea5 libc.so.6`abort + 295
    frame #2: 0x00007ffff5c21d79 libc.so.6`__assert_fail_base.cold.0 + 15
    frame #3: 0x00007ffff5c47426 libc.so.6`__assert_fail + 70
    frame #4: 0x000055555dc0923c clangd`clang::driver::InputInfo::getFilename(this=0x00007fffffff9398) const at InputInfo.h:84:5
    frame #5: 0x000055555dcd0d8d clangd`clang::driver::tools::SplitDebugName(JA=0x000055555f6c6a50, Args=0x000055555f6d0b80, Input=0x00007fffffff9678, Output=0x00007fffffff9398) at CommonArgs.cpp:1275:40
    frame #6: 0x000055555dc955a5 clangd`clang::driver::tools::Clang::ConstructJob(this=0x000055555f6c69d0, C=0x000055555f6c64a0, JA=0x000055555f6c6a50, Output=0x00007fffffff9398, Inputs=0x00007fffffff9668, Args=0x000055555f6d0b80, LinkingOutput=0x0000000000000000) const at Clang.cpp:5690:33
    frame #7: 0x000055555dbf6b54 clangd`clang::driver::Driver::BuildJobsForActionNoCache(this=0x00007fffffffb5e0, C=0x000055555f6c64a0, A=0x000055555f6c6a50, TC=0x000055555f6c4be0, BoundArch=(Data = 0x0000000000000000, Length = 0), AtTopLevel=true, MultipleArchs=false, LinkingOutput=0x0000000000000000, CachedResults=size=1, TargetDeviceOffloadKind=OFK_None) const at Driver.cpp:5618:10
    frame #8: 0x000055555dbf4ef0 clangd`clang::driver::Driver::BuildJobsForAction(this=0x00007fffffffb5e0, C=0x000055555f6c64a0, A=0x000055555f6c6a50, TC=0x000055555f6c4be0, BoundArch=(Data = 0x0000000000000000, Length = 0), AtTopLevel=true, MultipleArchs=false, LinkingOutput=0x0000000000000000, CachedResults=size=1, TargetDeviceOffloadKind=OFK_None) const at Driver.cpp:5306:26
    frame #9: 0x000055555dbeb590 clangd`clang::driver::Driver::BuildJobs(this=0x00007fffffffb5e0, C=0x000055555f6c64a0) const at Driver.cpp:4844:5
    frame #10: 0x000055555dbe6b0f clangd`clang::driver::Driver::BuildCompilation(this=0x00007fffffffb5e0, ArgList=ArrayRef<const char *> @ 0x00007fffffffb268) at Driver.cpp:1496:3
    frame #11: 0x000055555b0cc0d9 clangd`clang::createInvocation(ArgList=ArrayRef<const char *> @ 0x00007fffffffbb38, Opts=CreateInvocationOptions @ 0x00007fffffffbb90) at CreateInvocationFromCommandLine.cpp:53:52
    frame #12: 0x000055555b378e7b clangd`clang::clangd::buildCompilerInvocation(Inputs=0x00007fffffffca58, D=0x00007fffffffc158, CC1Args=size=0) at Compiler.cpp:116:44
    frame #13: 0x000055555895a6c8 clangd`clang::clangd::(anonymous namespace)::Checker::buildInvocation(this=0x00007fffffffc760, TFS=0x00007fffffffe570, Contents= Has Value=false ) at Check.cpp:212:9
    frame #14: 0x0000555558959cec clangd`clang::clangd::check(File=(Data = "build/test.cpp", Length = 64), TFS=0x00007fffffffe570, Opts=0x00007fffffffe600) at Check.cpp:486:34
    frame #15: 0x000055555892164a clangd`main(argc=4, argv=0x00007fffffffecd8) at ClangdMain.cpp:993:12
    frame #16: 0x00007ffff5c3ad85 libc.so.6`__libc_start_main + 229
    frame #17: 0x00005555585bbe9e clangd`_start + 46
```

Test Plan: ninja ClangDriverTests && tools/clang/unittests/Driver/ClangDriverTests

Differential Revision: https://reviews.llvm.org/D154602
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants