This is an archive of the discontinued LLVM Phabricator instance.

Differential D83478

[OPENMP]Fix compiler crash for target data directive without actual target codegen.
ClosedPublic

Authored by ABataev on Jul 9 2020, 5:57 AM.

Details

Reviewers
jdoerfert
Commits
rG7075c056e91b: [OPENMP]Fix compiler crash for target data directive without actual target…
Summary

Need to privatize addresses of the captured variables when trying to
emit the body of the target data directive in no target codegen mode.

Diff Detail

Event Timeline

ABataev created this revision.Jul 9 2020, 5:57 AM
Herald added a project: Restricted Project. · View Herald TranscriptJul 9 2020, 5:57 AM
Herald added subscribers: sstefan1, guansong, yaxunl. · View Herald Transcript
Harbormaster failed remote builds in B63578: Diff 276720!Jul 9 2020, 6:40 AM
jdoerfert accepted this revision.Jul 10 2020, 5:04 PM

LGTM

This revision is now accepted and ready to land.Jul 10 2020, 5:04 PM
Closed by commit rG7075c056e91b: [OPENMP]Fix compiler crash for target data directive without actual target… (authored by ABataev). · Explain WhyJul 13 2020, 7:55 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
CodeGen/
1 line
test/
OpenMP/
19 lines

Diff 277438

clang/lib/CodeGen/CGStmtOpenMP.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 6,071 Lines • ▼ Show 20 Lines auto &&PrivCodeGen = [&S, &Info, &PrivatizeDevicePointers,
CGF.EmitOMPUseDevicePtrClause(*C, PrivateScope, CGF.EmitOMPUseDevicePtrClause(*C, PrivateScope,
Info.CaptureDeviceAddrMap); Info.CaptureDeviceAddrMap);
for (const auto *C : S.getClausesOfKind<OMPUseDeviceAddrClause>()) for (const auto *C : S.getClausesOfKind<OMPUseDeviceAddrClause>())
CGF.EmitOMPUseDeviceAddrClause(*C, PrivateScope, CGF.EmitOMPUseDeviceAddrClause(*C, PrivateScope,
Info.CaptureDeviceAddrMap); Info.CaptureDeviceAddrMap);
(void)PrivateScope.Privatize(); (void)PrivateScope.Privatize();
RCG(CGF); RCG(CGF);
} else { } else {
OMPLexicalScope Scope(CGF, S, OMPD_unknown);
RCG(CGF); RCG(CGF);
} }
}; };
// Forward the provided action to the privatization codegen. // Forward the provided action to the privatization codegen.
RegionCodeGenTy PrivRCG(PrivCodeGen); RegionCodeGenTy PrivRCG(PrivCodeGen);
PrivRCG.setAction(Action); PrivRCG.setAction(Action);
▲ Show 20 Lines • Show All 548 Lines • Show Last 20 Lines

clang/test/OpenMP/target_data_codegen.cpp

Show First 20 Lines • Show All 485 Lines • ▼ Show 20 Lines
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}} // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
#ifdef CK6 #ifdef CK6
void test_close_modifier(int arg) { void test_close_modifier(int arg) {
// CK6: private unnamed_addr constant [1 x i64] [i64 1059] // CK6: private unnamed_addr constant [1 x i64] [i64 1059]
#pragma omp target data map(close,tofrom: arg) #pragma omp target data map(close,tofrom: arg)
{++arg;} {++arg;}
} }
#endif #endif
///==========================================================================///
// RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
// RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
// RUN: %clang_cc1 -DCK7 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY7 %s
// RUN: %clang_cc1 -DCK7 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY7 %s
// SIMD-ONLY7-NOT: {{__kmpc|__tgt}}
#ifdef CK7
// CK7: test_device_ptr_addr
void test_device_ptr_addr(int arg) {
int *p;
// CK7: add nsw i32
// CK7: add nsw i32
#pragma omp target data use_device_ptr(p) use_device_addr(arg)
{ ++arg, ++(*p); }
}
#endif
#endif #endif