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

[sdmitriev]

Summary

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.

[mgrang]

Please add some unit tests.

lib/CodeGen/CGOpenMPRuntime.cpp

3830

Please use llvm::sort instead of std::sort. See https://llvm.org/docs/CodingStandards.html#beware-of-non-deterministic-sorting-order-of-equal-elements.

[sdmitriev]

Replaced std::sort with llvm::sort. Added a test for offload target registration code for two offload targets.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3825

Preallocate the memory for all elements at construction, you know the number of elements.

[sdmitriev]

lib/CodeGen/CGOpenMPRuntime.cpp

3825

Agree.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3828–3829

Use something like this:

llvm::copy(Devices, std::next(RegFnNameParts.begin(), 2));

3830

Also, use std::next(RegFnNameParts.begin(), 2)

[ABataev]

LG

lib/CodeGen/CGOpenMPRuntime.cpp

CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
   llvm::Function *RegFn;
   {
     CodeGenFunction CGF(CGM);
     // Disable debug info for global (de-)initializer because they are not part
     // of some particular construct.
     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);
     // Create a variable to drive the registration and unregistration of the
     // descriptor, so we can reuse the logic that emits Ctors and Dtors.
     ImplicitParamDecl RegUnregVar(C, C.getTranslationUnitDecl(),
                                   SourceLocation(), nullptr, C.CharTy,
                                   ImplicitParamDecl::Other);

test/OpenMP/openmp_offload_registration.cpp

+// 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]]*)
+