[OpenMP] Reorganize code to allow specialized code generation for different devices.

[sfantao]

Summary

Different devices may in some cases require different code generation schemes in order to implement OpenMP. This is required not only for performance reasons, but also because it may not be possible to have the current (default) implementation working for these devices. E.g. GPU's cannot implement the same scheme a target such as powerpc or x86b would use, in the sense that it does not have the ability to fork threads, instead all the threads are always executing and need to be managed by the implementation.

This patch proposes a reorganization of the code in the OpenMP code generation to pave the way to have specialized implementation of OpenMP support. More than a "real" patch this is more a request for comments in order to understand if what is proposed is acceptable or if there are better/easier ways to do it.

In this patch part of the common OpenMP codegen infrastructure is moved to a new file under a new namespace (CGOpenMPCommon) so it can be shared between the default implementation and the specialized one. When CGOpenMPRuntime is created, an attempt to select a specialized implementation is done.

In the patch a specialization for nvptx targets is done which currently checks if the target is an OpenMP device and trap if it is not.

Let me know comments suggestions you may have.

[ABataev]

lib/CodeGen/CGOpenMPRuntimeCommon.h

1 ↗

(On Diff #46566)

I don't think we need this new file and new namespace. If some (currently) internal classes are needed, they must be exposed via CGOpenMPRuntime class. Derived class will be able to use these classes.
Also, do not expose classes if they are not required right now.

267–268 ↗

(On Diff #46566)

I think it is better to make this function a member of CodeGenFunction.

270–279 ↗

(On Diff #46566)

If you need this one, make it a virtual member of CGOpenMPRuntime

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

25

I think it must be checked during creation of NVPTX specific OpenMPRuntime instance.

[sfantao]

Clean up changes that are not required now.

Use CGOpenMPRuntime to contain everything that requires sharing.

Create diagnostic to notify user about unsupported OpenMP targets.

[sfantao]

Hi Alexey,

Thanks again for the review!

lib/CodeGen/CGOpenMPRuntimeCommon.h

1 ↗

(On Diff #46566)

Alright. I will add what is required as protected members of CGOpenMPRuntime. For now I didn't do that for any of the classes given that we are not doing any actual codegen.

267–268 ↗

(On Diff #46566)

Ok, once we start moving things to CGOpenMPRntime runtime, we add this to CodeGenFunction accordingly.

270–279 ↗

(On Diff #46566)

Ok. I'll do that once we post the codegen patches.

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

25

Ok, I did that. Also I added a new diagnostic in the compiler invocation so that the user gets a message instead of a stack dump. Let me know if you agree.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

340–341 ↗

(On Diff #46825)

I think it is better to reorganize this code like that:

switch (CGM.getTarget().getTriple().getArch()) {
  case llvm::Triple::nvptx:
  case llvm::Triple::nvptx64:
    assert(CGM.getLangOpts().OpenMPIsDevice &&
           "OpenMP NVPTX is only prepared to deal with device code.");
    return new CGOpenMPRuntimeNVPTX(CGM);
  default:
    break;
}
return new CGOpenMPRuntime(CGM);

lib/CodeGen/CGOpenMPRuntime.h

552–555 ↗

(On Diff #46825)

Do you really need this function? Currently, I don't see a point in adding this platform-specific thing to (mostly) common interface

[sfantao]

Hi Alexey,

Thanks for the feedback!

lib/CodeGen/CGOpenMPRuntime.cpp

340–341 ↗

(On Diff #46825)

Ok, will do.

lib/CodeGen/CGOpenMPRuntime.h

552–555 ↗

(On Diff #46825)

I am using that function to avoid exposing the constructor of CGOpenMPRuntimeNVPTX in CGOpenMPRuntime.h. Do you prefer me to do that instead?

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

552–555 ↗

(On Diff #46825)

Why do you need to expose constructor of CGOpenMPRuntimeNVPTX in CGOpenMPRuntime.h? It must be exposed only in CGOpenMPRuntime.cpp. Also, it would be a good idea to remove CGOpenMPRuntime::create() and move the whole logic to CodeGenModule::createOpenMPRuntime()

[sfantao]

lib/CodeGen/CGOpenMPRuntime.h

552–555 ↗

(On Diff #46825)

Ok I can move things to CodeGenModule.

I am afraid I am not understanding what you want me to do with CGOpenMPRuntimeNVPTX constructor. We have CGOpenMPRuntimeNVPTX in CGOpenMPRuntimeNVPTX.cpp, so unless I declare it in CGOpenMPRuntime.h (or some other header file) or have a function declared somewhere and defined in CGOpenMPRuntimeNVPTX.cpp, I can't call the Ctor.

So do you want me do create two functions in the CodeGen namespace (say createDefaultOpenMPRuntime() and createNVPTXOpenMPRuntime) and use them in CodeGenModule after checking which target is being used?

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

552–555 ↗

(On Diff #46825)

Oh, I missed that you did not add header file. I think you need to add a new one for new runtime support library.

[sfantao]

Create CGOpenMPRuntimeNVPTX header file and mike specialization selection in CodeGenModule.

[ABataev]

lib/CodeGen/CodeGenModule.cpp

218

I think this must be removed

lib/Frontend/CompilerInvocation.cpp

1837

OpenMPIsDevice is always false here. Do we really need it here?

[sfantao]

Remove return statement and use 'false' instead of relying on OpenMPIsDevice.

lib/CodeGen/CodeGenModule.cpp

218

Ok, done.

lib/Frontend/CompilerInvocation.cpp

1837

Ok, hardcoded it to false.

[sfantao]

Add comment.

[ABataev]

lib/Frontend/CompilerInvocation.cpp

1837

No, I mean do we really need this argument if this is a constant? I think for now we should just throw it away and edit error message so that it does not expect the second argument,

[sfantao]

Change diagnostic for invalid host target.

[sfantao]

lib/Frontend/CompilerInvocation.cpp

1837

Ok, I did as you say. thanks!

[ABataev]

LG

include/clang/Basic/DiagnosticDriverKinds.td

 def err_drv_emit_llvm_link : Error<
    "-emit-llvm cannot be used when linking">;
 def err_drv_optimization_remark_pattern : Error<
   "%0 in '%1'">;
 def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
 def err_drv_invalid_omp_target : Error<"OpenMP target is invalid: '%0'">;
 def err_drv_omp_host_ir_file_not_found : Error<
   "The provided host compiler IR file '%0' is required to generate code for OpenMP target regions but cannot be found.">;
+def err_drv_omp_host_target_not_supported : Error<
+  "The target '%0' is not a supported OpenMP host target.">;
   
 def warn_O4_is_O3 : Warning<"-O4 is equivalent to -O3">, InGroup<Deprecated>;
 def warn_drv_lto_libpath : Warning<"libLTO.dylib relative to clang installed dir not found; using 'ld' default search path instead">,
   InGroup<LibLTO>;
 def warn_drv_optimization_value : Warning<"optimization level '%0' is not supported; using '%1%2' instead">,
   InGroup<InvalidCommandLineArgument>;
 def warn_ignored_gcc_optimization : Warning<"optimization flag '%0' is not supported">,
   InGroup<IgnoredOptimizationArgument>;
 def warn_c_kext : Warning<

lib/CodeGen/CGOpenMPRuntimeNVPTX.h

+//===----- CGOpenMPRuntimeNVPTX.h - Interface to OpenMP NVPTX Runtimes ----===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides a class for OpenMP runtime code generation specialized to NVPTX
+// targets.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H
+#define LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H
+
+#include "CGOpenMPRuntime.h"
+
+namespace clang {
+namespace CodeGen {
+
+class CGOpenMPRuntimeNVPTX : public CGOpenMPRuntime {
+public:
+  explicit CGOpenMPRuntimeNVPTX(CodeGenModule &CGM);
+};
+
+} // CodeGen namespace.
+} // clang namespace.
+
+#endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

+//===---- CGOpenMPRuntimeNVPTX.cpp - Interface to OpenMP NVPTX Runtimes ---===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides a class for OpenMP runtime code generation specialized to NVPTX
+// targets.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CGOpenMPRuntimeNVPTX.h"
+
+using namespace clang;
+using namespace CodeGen;
+
+CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM)
+    : CGOpenMPRuntime(CGM) {}

lib/CodeGen/CMakeLists.txt

add_clang_library(clangCodeGen
   CGExprScalar.cpp
   CGLoopInfo.cpp
   CGObjC.cpp
   CGObjCGNU.cpp
   CGObjCMac.cpp
   CGObjCRuntime.cpp
   CGOpenCLRuntime.cpp
   CGOpenMPRuntime.cpp
+  CGOpenMPRuntimeNVPTX.cpp
   CGRecordLayoutBuilder.cpp
   CGStmt.cpp
   CGStmtOpenMP.cpp
   CGVTT.cpp
   CGVTables.cpp
   CodeGenABITypes.cpp
   CodeGenAction.cpp
   CodeGenFunction.cpp

lib/CodeGen/CodeGenModule.cpp

 #include "CGBlocks.h"
 #include "CGCUDARuntime.h"
 #include "CGCXXABI.h"
 #include "CGCall.h"
 #include "CGDebugInfo.h"
 #include "CGObjCRuntime.h"
 #include "CGOpenCLRuntime.h"
 #include "CGOpenMPRuntime.h"
+#include "CGOpenMPRuntimeNVPTX.h"
 #include "CodeGenFunction.h"
 #include "CodeGenPGO.h"
 #include "CodeGenTBAA.h"
 #include "CoverageMappingGen.h"
 #include "TargetInfo.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/CharUnits.h"
 #include "clang/AST/DeclCXX.h"
void CodeGenModule::createObjCRuntime() {
   llvm_unreachable("bad runtime kind");
 }
 
 void CodeGenModule::createOpenCLRuntime() {
   OpenCLRuntime = new CGOpenCLRuntime(*this);
 }
 
 void CodeGenModule::createOpenMPRuntime() {
+  // Select a specialized code generation class based on the target, if any.
+  // If it does not exist use the default implementation.
+  switch (getTarget().getTriple().getArch()) {
+
+  case llvm::Triple::nvptx:
+  case llvm::Triple::nvptx64:
+    assert(getLangOpts().OpenMPIsDevice &&
+           "OpenMP NVPTX is only prepared to deal with device code.");
+    OpenMPRuntime = new CGOpenMPRuntimeNVPTX(*this);
+    break;
+  default:
     OpenMPRuntime = new CGOpenMPRuntime(*this);
+    break;
+  }
 }
 
 void CodeGenModule::createCUDARuntime() {
   CUDARuntime = CreateNVCUDARuntime(*this);
 }
 
 void CodeGenModule::addReplacement(StringRef Name, llvm::Constant *C) {
   Replacements[Name] = C;
 }

lib/Frontend/CompilerInvocation.cpp

static Visibility parseVisibility(Arg *arg, ArgList &args,
   }
 
   diags.Report(diag::err_drv_invalid_value)
     << arg->getAsString(args) << value;
   return DefaultVisibility;
 }
 
 static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
+                          const TargetOptions &TargetOpts,
                           DiagnosticsEngine &Diags) {
   // FIXME: Cleanup per-file based stuff.
   LangStandard::Kind LangStd = LangStandard::lang_unspecified;
   if (const Arg *A = Args.getLastArg(OPT_std_EQ)) {
     LangStd = llvm::StringSwitch<LangStandard::Kind>(A->getValue())
 #define LANGSTANDARD(id, name, desc, features) \
       .Case(name, LangStandard::lang_##id)
 #include "clang/Frontend/LangStandards.def"
#include "clang/Frontend/LangStandards.def"
 
   // Check if -fopenmp is specified.
   Opts.OpenMP = Args.hasArg(options::OPT_fopenmp);
   Opts.OpenMPUseTLS =
       Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
   Opts.OpenMPIsDevice =
       Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_device);
 
+  // Provide diagnostic when a given target is not expected to be an OpenMP
+  // device or host.
+  if (Opts.OpenMP && !Opts.OpenMPIsDevice) {
+    llvm::Triple T(TargetOpts.Triple);
+    switch (T.getArch()) {
+    default:
+      break;
+    // Add unsupported host targets here:
+    case llvm::Triple::nvptx:
+    case llvm::Triple::nvptx64:
+      Diags.Report(clang::diag::err_drv_omp_host_target_not_supported)
+          << TargetOpts.Triple;
+      break;
+    }
+  }
+
   // Get the OpenMP target triples if any.
   if (Arg *A = Args.getLastArg(options::OPT_omptargets_EQ)) {
 
     for (unsigned i = 0; i < A->getNumValues(); ++i) {
       llvm::Triple TT(A->getValue(i));
 
       if (TT.getArch() == llvm::Triple::UnknownArch)
         Diags.Report(clang::diag::err_drv_invalid_omp_target) << A->getValue(i);
if (DashX == IK_AST || DashX == IK_LLVM_IR) {
     // PassManager in BackendUtil.cpp. They need to be initializd no matter
     // what the input type is.
     if (Args.hasArg(OPT_fobjc_arc))
       Res.getLangOpts()->ObjCAutoRefCount = 1;
     parseSanitizerKinds("-fsanitize=", Args.getAllArgValues(OPT_fsanitize_EQ),
                         Diags, Res.getLangOpts()->Sanitize);
   } else {
     // Other LangOpts are only initialzed when the input is not AST or LLVM IR.
-    ParseLangArgs(*Res.getLangOpts(), Args, DashX, Diags);
+    ParseLangArgs(*Res.getLangOpts(), Args, DashX, Res.getTargetOpts(), Diags);
     if (Res.getFrontendOpts().ProgramAction == frontend::RewriteObjC)
       Res.getLangOpts()->ObjCExceptions = 1;
   }
   // FIXME: ParsePreprocessorArgs uses the FileManager to read the contents of
   // PCH file and find the original header name. Remove the need to do that in
   // ParsePreprocessorArgs and remove the FileManager
   // parameters from the function and the "FileManager.h" #include.
   FileManager FileMgr(Res.getFileSystemOpts());

test/OpenMP/target_messages.cpp

 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 -o - %s
 // RUN: not %clang_cc1 -fopenmp -std=c++11 -omptargets=aaa-bbb-ccc-ddd -o - %s 2>&1 | FileCheck %s
 // CHECK: error: OpenMP target is invalid: 'aaa-bbb-ccc-ddd'
+// RUN: not %clang_cc1 -fopenmp -std=c++11 -triple nvptx64-nvidia-cuda -o - %s 2>&1 | FileCheck --check-prefix CHECK-UNSUPPORTED-HOST-TARGET %s
+// RUN: not %clang_cc1 -fopenmp -std=c++11 -triple nvptx-nvidia-cuda -o - %s 2>&1 | FileCheck --check-prefix CHECK-UNSUPPORTED-HOST-TARGET %s
+// CHECK-UNSUPPORTED-HOST-TARGET: error: The target '{{nvptx64-nvidia-cuda|nvptx-nvidia-cuda}}' is not a supported OpenMP host target.
 
 void foo() {
 }
 
 #pragma omp target // expected-error {{unexpected OpenMP directive '#pragma omp target'}}
 
 int main(int argc, char **argv) {
   #pragma omp target { // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}