This is an archive of the discontinued LLVM Phabricator instance.

Differential D50218

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

Authored by sdmitriev on Aug 2 2018, 8:16 PM.

Details

Reviewers
ABataev
hfinkel
Commits
rGbde9cf942b7c: [OpenMP] Encode offload target triples into comdat key for offload…
rL338916: [OpenMP] Encode offload target triples into comdat key for offload…
rC338916: [OpenMP] Encode offload target triples into comdat key for offload…
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.

Diff Detail

Event Timeline

sdmitriev created this revision.Aug 2 2018, 8:16 PM
Herald added subscribers: cfe-commits, mgrang, guansong. · View Herald TranscriptAug 2 2018, 8:16 PM
mgrang added a comment.Aug 2 2018, 8:55 PM

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 updated this revision to Diff 159028.Aug 3 2018, 9:36 AM

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

ABataev added inline comments.Aug 3 2018, 10:37 AM
lib/CodeGen/CGOpenMPRuntime.cpp
3825

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

sdmitriev added inline comments.Aug 3 2018, 10:52 AM
lib/CodeGen/CGOpenMPRuntime.cpp
3825

Agree.

sdmitriev updated this revision to Diff 159058.Aug 3 2018, 11:15 AM
ABataev added inline comments.Aug 3 2018, 11:28 AM
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)

sdmitriev updated this revision to Diff 159074.Aug 3 2018, 1:09 PM
ABataev accepted this revision.Aug 3 2018, 1:12 PM

LG

This revision is now accepted and ready to land.Aug 3 2018, 1:12 PM
Closed by commit rC338916: [OpenMP] Encode offload target triples into comdat key for offload… (authored by sdmitriev). · Explain WhyAug 3 2018, 1:20 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
lib/
CodeGen/
12 lines
test/
OpenMP/
49 lines

Diff 159028

lib/CodeGen/CGOpenMPRuntime.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 3,812 Lines • ▼ Show 20 LinesCGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
llvm::Function *RegFn; llvm::Function *RegFn;
{ {
CodeGenFunction CGF(CGM); CodeGenFunction CGF(CGM);
// Disable debug info for global (de-)initializer because they are not part // Disable debug info for global (de-)initializer because they are not part
// of some particular construct. // of some particular construct.
CGF.disableDebugInfo(); CGF.disableDebugInfo();
const auto &FI = CGM.getTypes().arrangeNullaryFunction(); const auto &FI = CGM.getTypes().arrangeNullaryFunction();
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI); 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;
ABataevUnsubmitted
Not Done
ReplyInline Actions

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

ABataev: Preallocate the memory for all elements at construction, you know the number of elements.
sdmitrievAuthorUnsubmitted
Not Done
ReplyInline Actions

Agree.

sdmitriev: Agree.
RegFnNameParts.push_back("omp_offloading");
RegFnNameParts.push_back("descriptor_reg");
for (const auto &Device : Devices)
RegFnNameParts.push_back(Device.getTriple());
ABataevUnsubmitted
Not Done
ReplyInline Actions

Use something like this:

llvm::copy(Devices, std::next(RegFnNameParts.begin(), 2));
ABataev: Use something like this: ``` llvm::copy(Devices, std::next(RegFnNameParts.begin(), 2)); ```
llvm::sort(RegFnNameParts.begin() + 2, RegFnNameParts.end());
mgrangUnsubmitted
Not Done
ReplyInline Actions

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

mgrang: Please use llvm::sort instead of std::sort. See https://llvm.org/docs/CodingStandards.
ABataevUnsubmitted
Not Done
ReplyInline Actions

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

ABataev: Also, use `std::next(RegFnNameParts.begin(), 2)`
std::string Descriptor = getName(RegFnNameParts);
RegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, Descriptor, FI); RegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, Descriptor, FI);
CGF.StartFunction(GlobalDecl(), C.VoidTy, RegFn, FI, FunctionArgList()); CGF.StartFunction(GlobalDecl(), C.VoidTy, RegFn, FI, FunctionArgList());
CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_register_lib), Desc); CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_register_lib), Desc);
// Create a variable to drive the registration and unregistration of the // Create a variable to drive the registration and unregistration of the
// descriptor, so we can reuse the logic that emits Ctors and Dtors. // descriptor, so we can reuse the logic that emits Ctors and Dtors.
ImplicitParamDecl RegUnregVar(C, C.getTranslationUnitDecl(), ImplicitParamDecl RegUnregVar(C, C.getTranslationUnitDecl(),
SourceLocation(), nullptr, C.CharTy, SourceLocation(), nullptr, C.CharTy,
ImplicitParamDecl::Other); ImplicitParamDecl::Other);
▲ Show 20 Lines • Show All 5,364 Lines • Show Last 20 Lines

test/OpenMP/openmp_offload_registration.cpp

PropertyOld ValueNew Value
svn:eol-stylenullnative
svn:keywordsnullAuthor Date Id Rev URL
svn:mime-typenulltext/plain
// 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]]*)