This is an archive of the discontinued LLVM Phabricator instance.

[OpenMP] Fix map clause for unused var: don't ignore it
ClosedPublic

Authored by jdenny on Jul 15 2020, 6:06 PM.

Details

Summary

For example, without this patch:

$ cat test.c
int main() {
  int x[3];
  #pragma omp target map(tofrom:x[0:3])
#ifdef USE 
  x[0] = 1
#endif
  ;
  return 0;
}
$ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -S -emit-llvm test.c
$ grep '^@.offload_maptypes' test.ll
$ echo $?
1
$ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -S -emit-llvm test.c \
        -DUSE
$ grep '^@.offload_maptypes' test.ll
@.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35] 

With this patch, both greps produce the same result.

Diff Detail

Event Timeline

jdenny created this revision.Jul 15 2020, 6:06 PM

I would add checks for mapping of declare target to/link vars and checked if they work in runtime.

clang/lib/CodeGen/CGOpenMPRuntime.cpp
9468–9469

I would check that we capture a variable. We may capture VLA here as well.

9470–9471

No need to insert nullptr here, it is not used later.

ABataev added inline comments.Jul 16 2020, 7:25 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7946–7947

Use llvm::DenseSet<CanonicalDeclPtr<const ValueDecl>> &SkipVarSet.

9439

llvm::DenseSet<CanonicalDeclPtr<const ValueDecl>> MappedVarSet;

9469

No need to get canonical decl here for llvm::DenseSet<CanonicalDeclPtr<const ValueDecl>>

jdenny updated this revision to Diff 278546.Jul 16 2020, 11:13 AM
jdenny marked 7 inline comments as done.
jdenny edited the summary of this revision. (Show Details)

Use CanonicalDeclPtr, as suggested.

jdenny marked an inline comment as done.Jul 16 2020, 11:14 AM
jdenny added inline comments.
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7946–7947

Thanks for the suggestion.

CanonicalDeclPtr<const ValueDecl> appears to be unusable. The constructor (in Redeclarable.h) tries to initialize the const ValueDecl *Ptr member with the Decl * returned by getCanonicalDecl, producing a type mismatch.

I went with CanonicalDeclPtr<const Decl> instead.

9468–9469

I would check that we capture a variable. We may capture VLA here as well.

The if on line 9454 above checks for VLAs. This code is in the else.

Similarly, the existing implementation for generateInfoForCapture, which is called within this same else, doesn't protect against VLAs. It just checks capturesThis as I'm doing there.

Have I misunderstood?

9470–9471

Without this nulltpr insertion, many tests fail because of duplicate map entries. Independently of my patch, nulltptr is used to indicate this capture (see the generateInfoForCapture implementation) .

For example:

struct S {
  int i;
  void foo() {
    #pragma omp target map(i)
    i = 5;
  }
};

We should have:

@.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710659]

Without the nullptr insertion, we have:

@.offload_maptypes = private unnamed_addr constant [4 x i64] [i64 32, i64 281474976710659, i64 32, i64 844424930131971]

I would add checks for mapping of declare target to/link vars and checked if they work in runtime.

There are existing codegen tests for that, and they don't seem to be affected by this patch. This patch only addresses the case where a map clause is specified for an unused variable. Is there another behavior this patch might impact?

I would add checks for mapping of declare target to/link vars and checked if they work in runtime.

There are existing codegen tests for that, and they don't seem to be affected by this patch. This patch only addresses the case where a map clause is specified for an unused variable. Is there another behavior this patch might impact?

What I mean is the explicit mapping for declare target to/link variables. The variable is marked as declare to and then mapped. Do we have the tests for something like this?

clang/lib/CodeGen/CGOpenMPRuntime.cpp
9468–9469

Ah, yes, here we processed VLAs already.

9470–9471

This is strange, because you don't check for nullptr. You only check for ValueDecls, but not capture of this.

jdenny updated this revision to Diff 278605.Jul 16 2020, 2:45 PM
jdenny marked 3 inline comments as done.

Extended test coverage to this pointer when variable is unused.

I would add checks for mapping of declare target to/link vars and checked if they work in runtime.

There are existing codegen tests for that, and they don't seem to be affected by this patch. This patch only addresses the case where a map clause is specified for an unused variable. Is there another behavior this patch might impact?

What I mean is the explicit mapping for declare target to/link variables. The variable is marked as declare to and then mapped. Do we have the tests for something like this?

I didn't find tests for the case where a variable has a declare target and a map clause. I can add some to show that this patch doesn't change the generated map types.

This patch doesn't affect the example in that bug report as far as I can tell.

clang/lib/CodeGen/CGOpenMPRuntime.cpp
9470–9471

I'm not sure what code you're referring to when you say "you check".

Both with and without my patch, my understanding is that nullptr is a special key that means this. I'm depending on that to avoid generating map entries twice for this.

My understanding is based on the way generateInfoForCapture works. If Cap->capturesThis(), then VD = nullptr. That VD is then used by C->decl_component_lists(VD) to find entries for this in map clauses.

Unless I'm misreading it, the code that sets nullptr for this in decl components is checkMappableExpressionList in SemaOpenMP.cpp. The last few lines of that function have a comment to this effect. (That comment would probably be more useful in a header file somewhere.)

ABataev added inline comments.Jul 17 2020, 7:48 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
9470–9471

MappedVarSet is a new variable, right? It is used only in 2 places: here, where you add elements to this set, and in generateAllInfo where you have a check:

if (SkipVarSet.count(VD))
            return;

I don't see other places where it is used. And I don't see places where you check for the presence of nullptr in this set. That's why I think you don't need to add it, if you don't check for its presence later.

jdenny marked an inline comment as done.Jul 17 2020, 8:12 AM
jdenny added inline comments.
clang/lib/CodeGen/CGOpenMPRuntime.cpp
9470–9471

MappedVarSet is a new variable, right?

Right.

It is used only in 2 places: here, where you add elements to this set, and in generateAllInfo where you have a check:

if (SkipVarSet.count(VD))
            return;

I don't see other places where it is used.

That's all.

And I don't see places where you check for the presence of nullptr in this set. That's why I think you don't need to add it, if you don't check for its presence later.

The SkipVarSet.count(VD) you mentioned checks for presence of any key in order to avoid creating duplicate map entries. nullptr is one such key. It happens to indicate this.

For comparison, look inside generateInfoForCapture where the following code establishes that nullptr is the key for this:

const ValueDecl *VD = Cap->capturesThis()
                          ? nullptr
                          : Cap->getCapturedVar()->getCanonicalDecl();

After that, does generateInfoForCapture ever check for VD == nullptr? I don't see where it does. But it does use VD, which might be nulltpr, as a decl components key in the following code:

for (const auto L : C->decl_component_lists(VD)) {

Likewise, my code is also using a VD that might be nullptr as a key in SkipVarSet. I don't have to special case nullptr at this point. It's just another key.

ABataev accepted this revision.Jul 17 2020, 8:18 AM

LG

clang/lib/CodeGen/CGOpenMPRuntime.cpp
9470–9471

Ah, I see, I missed the previous line.

This revision is now accepted and ready to land.Jul 17 2020, 8:18 AM
jdenny marked 3 inline comments as done.Jul 17 2020, 8:24 AM

Thanks for the review. Will try to push soon.

jdenny updated this revision to Diff 278933.Jul 17 2020, 4:49 PM

Rebased onto a more recent master.

This revision was automatically updated to reflect the committed changes.
ABataev added a comment.EditedJul 20 2020, 8:24 AM

Looks like it breaks unified_shared_memory/close_enter_exit.c test. Could you check this, please?

Looks like it breaks unified_shared_memory/close_enter_exit.c test. Could you check this, please?

It doesn't fail for me, and I didn't receive a buildbot notification. Do you have a link?

Looks like it breaks unified_shared_memory/close_enter_exit.c test. Could you check this, please?

It doesn't fail for me, and I didn't receive a buildbot notification. Do you have a link?

I saw it failed it in local build, but probably it was caused by a misconfiguration, sorry.