Runtime crash with use of requires unified_shared_memory clause

[jyu2-git]

The problem is there is an extra indirection involved when requires shared memory is used. It seems like there is a missing load from the reference here.

 if (const auto *VD =
              dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) {
        if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
                OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
          if ((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
              (*Res == OMPDeclareTargetDeclAttr::MT_To &&
               CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) {
            RequiresReference = true;
            BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
          }
        }

 if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
            !VD || VD->hasLocalStorage())
          BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
        else
          FirstPointerInComplexData = true;

Reproduce:

$ clang -fopenmp -fopenmp-targets=x86_64 x.c
clang-15: /localdisk2/jyu2/git/xmain1/llvm/clang/lib/CodeGen/Address.h:103: clang::CodeGen::Address::Address(llvm::Value *, llvm::Type *, clang::CharUnits): Assertion `llvm::cast<llvm::PointerType>(Pointer->getType()) ->isOpaqueOrPointeeTypeMatches(ElementType) && "Incorrect pointer element type"' failed.
PLEASE append the compiler options "-save-temps -v", rebuild the application to to get the full command which is failing and submit a bug report to https://software.intel.com/en-us/support/priority-support which includes the failing command, input files for the command and the crash backtrace (if any).

There is the test case x.c:

#include <stdio.h>
#include <malloc.h>

#pragma omp requires unified_shared_memory
struct NODE
{
  int a;
  int b;
};

struct NODE *Nodes_base;
#pragma omp declare target to(Nodes_base)
void foo()
{
   Nodes_base = malloc(8);
   Nodes_base[0].a = 1;
   Nodes_base[0].b = 2;
   printf("berfor a = %d %d\n", Nodes_base[0].a, Nodes_base[0].b);
   #pragma omp target map(tofrom:Nodes_base[0:1])
     {
       Nodes_base[0].a = 3;
       Nodes_base[0].b = 4;
       printf("a = %d %d\n", Nodes_base[0].a, Nodes_base[0].b);
     }
   printf("after a = %d %d\n", Nodes_base[0].a, Nodes_base[0].b);
}
int main() {
  foo();
  return 0;
}

[llvmbot]

@llvm/issue-subscribers-openmp

[jhuber6]

I'm unable to reproduce this crash with the provided source and options, see https://godbolt.org/z/4GbjMarfh. Is this a problem in LLVM/Clang? Or is it only present in Intel's fork.

[jyu2-git]

I'm unable to reproduce this crash with the provided source and options, see https://godbolt.org/z/4GbjMarfh. Is this a problem in LLVM/Clang? Or is it only present in Intel's fork.

Sorry, the assertion only happen with -no-opaque-pointers. With opaqque-pointers turn on by default, the assertion can no longer reproduced. But the test still run failed with: cmplrllvm-39871>./a.out berfor a = 1 2 a.out: /localdisk2/jyu2/git/xmain1/llvm/openmp/libomptarget/src/omptarget.cpp:172: int initLibrary(DeviceTy &): Assertion `CurrDeviceEntry->size == CurrHostEntry->size && "data size mismatch"' failed. Aborted (core dumped)

[shiltian]

Is it still an issue?

[jyu2-git]

Yes, could you reproduce the problem with my test. I still see the runtime as: berfor a = 1 2 Segmentation fault (core dumped)

I have a fix in our source base.

[shiltian]

@jyu2-git can you upstream it and have Fix #57855. at the end of your commit message such that this issue can be closed automatically? It would be great if you can make it before the branching of LLVM 16. Thanks!

I can see the problem is we skip the processing if the argument passed to libomptarget is not a kernel argument.