r338916 - [OpenMP] Encode offload target triples into comdat key for offload initialization code

Sergey Dmitriev via cfe-commits cfe-commits at lists.llvm.org
Fri Aug 3 13:19:28 PDT 2018


Author: sdmitriev
Date: Fri Aug  3 13:19:28 2018
New Revision: 338916

URL: http://llvm.org/viewvc/llvm-project?rev=338916&view=rev
Log:
[OpenMP] Encode offload target triples into comdat key for offload initialization code

Encoding offload target triples onto comdat group key for offload initialization
code guarantees that it will be executed once per each unique combination of
offload targets.

Differential Revision: https://reviews.llvm.org/D50218

Added:
    cfe/trunk/test/OpenMP/openmp_offload_registration.cpp   (with props)
Modified:
    cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp

Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=338916&r1=338915&r2=338916&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Fri Aug  3 13:19:28 2018
@@ -3818,7 +3818,19 @@ CGOpenMPRuntime::createOffloadingBinaryD
     CGF.disableDebugInfo();
     const auto &FI = CGM.getTypes().arrangeNullaryFunction();
     llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
-    std::string Descriptor = getName({"omp_offloading", "descriptor_reg"});
+
+    // Encode offload target triples into the registration function name. It
+    // will serve as a comdat key for the registration/unregistration code for
+    // this particular combination of offloading targets.
+    SmallVector<StringRef, 4U> RegFnNameParts(Devices.size() + 2U);
+    RegFnNameParts[0] = "omp_offloading";
+    RegFnNameParts[1] = "descriptor_reg";
+    llvm::transform(Devices, std::next(RegFnNameParts.begin(), 2),
+                    [](const llvm::Triple &T) -> const std::string& {
+                      return T.getTriple();
+                    });
+    llvm::sort(std::next(RegFnNameParts.begin(), 2), RegFnNameParts.end());
+    std::string Descriptor = getName(RegFnNameParts);
     RegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, Descriptor, FI);
     CGF.StartFunction(GlobalDecl(), C.VoidTy, RegFn, FI, FunctionArgList());
     CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_register_lib), Desc);

Added: cfe/trunk/test/OpenMP/openmp_offload_registration.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/openmp_offload_registration.cpp?rev=338916&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/openmp_offload_registration.cpp (added)
+++ cfe/trunk/test/OpenMP/openmp_offload_registration.cpp Fri Aug  3 13:19:28 2018
@@ -0,0 +1,49 @@
+// Test for offload registration code for two targets
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu -fopenmp-targets=x86_64-pc-linux-gnu,powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+void foo() {
+#pragma omp target
+  {}
+}
+
+// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
+// CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
+// CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
+
+// Comdat key for the offload registration code. Should have sorted offload
+// target triples encoded into the name.
+// CHECK-DAG: $[[REGFN:\.omp_offloading\..+\.powerpc64le-ibm-linux-gnu\.x86_64-pc-linux-gnu+]] = comdat any
+
+// Check if offloading descriptor is created.
+// CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
+// CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
+// CHECK: [[DEV1BEGIN:@.+]] = extern_weak constant i8
+// CHECK: [[DEV1END:@.+]] = extern_weak constant i8
+// CHECK: [[DEV2BEGIN:@.+]] = extern_weak constant i8
+// CHECK: [[DEV2END:@.+]] = extern_weak constant i8
+// CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [2 x [[DEVTY]]] [{{.+}} { i8* [[DEV1BEGIN]], i8* [[DEV1END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, {{.+}} { i8* [[DEV2BEGIN]], i8* [[DEV2END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
+// CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 2, [[DEVTY]]* getelementptr inbounds ([2 x [[DEVTY]]], [2 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]])
+
+// Check target registration is registered as a Ctor.
+// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
+
+// Check presence of foo() and the outlined target region
+// CHECK: define void [[FOO:@.+]]()
+// CHECK: define internal void [[OUTLINEDTARGET:@.+]]() 
+
+// Check registration and unregistration code.
+
+// CHECK:     define internal void @[[UNREGFN:.+]](i8*)
+// CHECK-SAME: comdat($[[REGFN]]) {
+// CHECK:     call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
+// CHECK:     ret void
+// CHECK:     declare i32 @__tgt_unregister_lib([[DSCTY]]*)
+
+// CHECK:     define linkonce hidden void @[[REGFN]]()
+// CHECK-SAME: comdat {
+// CHECK:     call i32 @__tgt_register_lib([[DSCTY]]* [[DESC]])
+// CHECK:     call i32 @__cxa_atexit(void (i8*)* @[[UNREGFN]], i8* bitcast ([[DSCTY]]* [[DESC]] to i8*),
+// CHECK:     ret void
+// CHECK:     declare i32 @__tgt_register_lib([[DSCTY]]*)
+

Propchange: cfe/trunk/test/OpenMP/openmp_offload_registration.cpp
------------------------------------------------------------------------------
    svn:eol-style = native

Propchange: cfe/trunk/test/OpenMP/openmp_offload_registration.cpp
------------------------------------------------------------------------------
    svn:keywords = Author Date Id Rev URL

Propchange: cfe/trunk/test/OpenMP/openmp_offload_registration.cpp
------------------------------------------------------------------------------
    svn:mime-type = text/plain




More information about the cfe-commits mailing list