[WIP][OpenMP] Propagate 'declare target' information to device

[Hahnfeld]

Summary

Emit IR metadata with list of all functions found during host CodeGen.

[Hahnfeld]

This doesn't work the way I hoped :-(

The problem is that implicit declare target isn't determined during Sema and not present in AST. Instead this is handled in CodeGen (using addDeferredDeclToEmit if I understand rL327636 correctly).
So unless we reconsider something like D38798 the host doesn't know which functions are implicitly declare target and consequently can't propagate that information to the device compilation.

[ABataev]

D38798 does not work correctly and there is no stable Sema-based solution to do this absolutely correctly. We need to adapt the codegen so that it emits metadata for the implicit declare target functions.

Best regards,
Alexey Bataev

18 авг. 2018 г., в 10:26, Jonas Hahnfeld via Phabricator <[e-mail]> написал(а):

Hahnfeld added a comment.

This doesn't work the way I hoped :-(

The problem is that implicit declare target isn't determined during Sema and not present in AST. Instead this is handled in CodeGen (using addDeferredDeclToEmit if I understand https://reviews.llvm.org/rL327636 correctly).
So unless we reconsider something like https://reviews.llvm.org/D38798 the host doesn't know which functions are implicitly declare target and consequently can't propagate that information to the device compilation.

https://reviews.llvm.org/D50939

[Hahnfeld]

Okay, looks promising so far. I'll update this revision once ready for testing.

[Hahnfeld]

This revision correctly compiles the following two examples iff BREAK is not defined:

#include <stdint.h>

static void cpuid(int i, uint32_t *regs) {
  asm volatile("cpuid" : "=a" (regs[0]), "=b" (regs[1]), "=c" (regs[2]), "=d" (regs[3])
               : "a" (i), "c" (0));
}

void host() {
  uint32_t regs[4];
  cpuid(0, regs);
}

#pragma omp declare target
void target() {
#ifdef BREAK
  uint32_t regs[4];
  cpuid(0, regs);
#endif
}
#pragma omp end declare target

static void vla(int n) {
  int a[n];
}

void target() {
  vla(123);

#pragma omp target
{
#ifdef BREAK
  vla(123);
#endif
}
}

It doesn't work for C++ yet because the LLVM function names are already mangled, so we need to do the same when checking in Sema. I think the tricky thing will be templates: Because we don't know which specialization will eventually be emitted we might need to check all of them if a template is referenced in device code. Thoughts?

[ABataev]

In D50939#1211333, @Hahnfeld wrote:

It doesn't work for C++ yet because the LLVM function names are already mangled, so we need to do the same when checking in Sema. I think the tricky thing will be templates: Because we don't know which specialization will eventually be emitted we might need to check all of them if a template is referenced in device code. Thoughts?

I expected a lot of problems with C++ and I don't think we will be able to resolve this, especially with the templates. We need to try some other solutions.

[Hahnfeld]

In D50939#1211391, @ABataev wrote:

In D50939#1211333, @Hahnfeld wrote:

It doesn't work for C++ yet because the LLVM function names are already mangled, so we need to do the same when checking in Sema. I think the tricky thing will be templates: Because we don't know which specialization will eventually be emitted we might need to check all of them if a template is referenced in device code. Thoughts?

I expected a lot of problems with C++ and I don't think we will be able to resolve this, especially with the templates. We need to try some other solutions.

I tend to agree.

My next idea was to "abuse" -fdelayed-template-parsing:

#include <stdint.h>

template <typename T>
void cpuid(int i, uint32_t *regs) {
  asm volatile("cpuid" : "=a" (regs[0]), "=b" (regs[1]), "=c" (regs[2]), "=d" (regs[3])
               : "a" (i), "c" (0));
}

template <typename T>
void host() {
  uint32_t regs[4];
  cpuid<T>(0, regs);
}

#pragma omp declare target
void target() {
#ifdef BREAK
  uint32_t regs[4];
  cpuid<bool>(0, regs);
#endif
}
#pragma omp end declare target

Clang will delay checking templates (= functions outside of declare target) until they are instantiated during CodeGen. However we also need to find all target directives to outline the offload entries, but we can't know this without actually parsing (and checking) all functions. So I don't think this will work either.

lib/CodeGen/CodeGenModule.cpp

else if (const auto *MD = dyn_cast<CXXMethodDecl>(FD))
     if (MD->isVirtual())
       F->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
 
   // Don't emit entries for function declarations in the cross-DSO mode. This
   // is handled with better precision by the receiving DSO.
   if (!CodeGenOpts.SanitizeCfiCrossDso)
     CreateFunctionTypeMetadataForIcall(FD, F);
 
-  if (getLangOpts().OpenMP && FD->hasAttr<OMPDeclareSimdDeclAttr>())
+  if (getLangOpts().OpenMP) {
+    if (FD->hasAttr<OMPDeclareSimdDeclAttr>())
       getOpenMPRuntime().emitDeclareSimdFunction(FD, F);
+
+    if (!getLangOpts().OpenMPIsDevice && FD->hasAttr<OMPDeclareTargetDeclAttr>()) {
+      F->addFnAttr("omp-declare-target");
+    }
+  }
 }
 
 void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
   assert(!GV->isDeclaration() &&
          "Only globals with definition can force usage.");
   LLVMUsed.emplace_back(GV);
 }
 

lib/Sema/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
+  Core
   Support
   )
 
 if (MSVC)
   set_source_files_properties(SemaDeclAttr.cpp PROPERTIES COMPILE_FLAGS /bigobj)
   set_source_files_properties(SemaExpr.cpp PROPERTIES COMPILE_FLAGS /bigobj)
   set_source_files_properties(SemaTemplate.cpp PROPERTIES COMPILE_FLAGS /bigobj)
 endif()

lib/Sema/SemaDecl.cpp

 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/ParsedTemplate.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
 #include "clang/Sema/Template.h"
 #include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/Triple.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/Module.h"
 #include <algorithm>
 #include <cstring>
 #include <functional>
 
 using namespace clang;
 using namespace sema;
 
 Sema::DeclGroupPtrTy Sema::ConvertDeclToDeclGroup(Decl *Ptr, Decl *OwnedType) {
if (New->getDeclName() && AddToScope) {
     // declaration of this entity. Qualified lookup for such an entity should
     // only find this declaration if there is no visible declaration of it.
     bool AddToContext = !D.isRedeclaration() || !New->isLocalExternDecl();
     PushOnScopeChains(New, S, AddToContext);
     if (!AddToContext)
       CurContext->addHiddenDecl(New);
   }
 
-  if (isInOpenMPDeclareTargetContext())
+  if (getLangOpts().OpenMPIsDevice) {
+    if (auto *OpenMPHostIR = getASTContext().getOpenMPHostIR()) {
+      if (auto *FD = dyn_cast<FunctionDecl>(New)) {
+        // FIXME: This should be mangled.
+        StringRef Name = FD->getNameAsString();
+        auto *HostFunc = OpenMPHostIR->getFunction(Name);
+        if (HostFunc && HostFunc->hasFnAttribute("omp-declare-target")) {
+          checkDeclIsAllowedInOpenMPTarget(nullptr, New);
+        }
+      }
+      // FIXME: Handle other declarators.
+    }
+  } else if (isInOpenMPDeclareTargetContext())
     checkDeclIsAllowedInOpenMPTarget(nullptr, New);
 
   return New;
 }
 
 /// Helper method to turn variable array types into constant array
 /// types in certain situations which would otherwise be errors (for
 /// GCC compatibility).