[clang] f3c857f - [OPENMP50]Add basic codegen support for ancestor device modifier.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 18 15:03:56 PDT 2020
Author: Alexey Bataev
Date: 2020-03-18T17:53:18-04:00
New Revision: f3c857fae297a1420693a129a7bf39098b3f046a
URL: https://github.com/llvm/llvm-project/commit/f3c857fae297a1420693a129a7bf39098b3f046a
DIFF: https://github.com/llvm/llvm-project/commit/f3c857fae297a1420693a129a7bf39098b3f046a.diff
LOG: [OPENMP50]Add basic codegen support for ancestor device modifier.
If the ancestor device modifier is used and the value of the device
clause is evaluated to 1, the ancestor device shall be used for the
execution.
Since the reverse offloading is not supported yet, the target construct
execution is always initiated from the host, not from the device. So, if
the ancestor modifier is specified, just execute target region on the
host.
Added:
clang/test/OpenMP/target_device_codegen.cpp
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGStmtOpenMP.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index fa60221e8b59..c4e4578c4b04 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9483,7 +9483,7 @@ void CGOpenMPRuntime::emitTargetNumIterationsCall(
void CGOpenMPRuntime::emitTargetCall(
CodeGenFunction &CGF, const OMPExecutableDirective &D,
llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID, const Expr *IfCond,
- const Expr *Device,
+ llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device,
llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
const OMPLoopDirective &D)>
SizeEmitter) {
@@ -9507,6 +9507,16 @@ void CGOpenMPRuntime::emitTargetCall(
auto &&ThenGen = [this, Device, OutlinedFn, OutlinedFnID, &D, &InputInfo,
&MapTypesArray, &CS, RequiresOuterTask, &CapturedVars,
SizeEmitter](CodeGenFunction &CGF, PrePostActionTy &) {
+ if (Device.getInt() == OMPC_DEVICE_ancestor) {
+ // Reverse offloading is not supported, so just execute on the host.
+ if (RequiresOuterTask) {
+ CapturedVars.clear();
+ CGF.GenerateOpenMPCapturedVars(CS, CapturedVars);
+ }
+ emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedFn, CapturedVars);
+ return;
+ }
+
// On top of the arrays that were filled up, the target offloading call
// takes as arguments the device id as well as the host pointer. The host
// pointer is used by the runtime library to identify the current target
@@ -9521,9 +9531,13 @@ void CGOpenMPRuntime::emitTargetCall(
// Emit device ID if any.
llvm::Value *DeviceID;
- if (Device) {
- DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
- CGF.Int64Ty, /*isSigned=*/true);
+ if (Device.getPointer()) {
+ assert((Device.getInt() == OMPC_DEVICE_unknown ||
+ Device.getInt() == OMPC_DEVICE_device_num) &&
+ "Expected device_num modifier.");
+ llvm::Value *DevVal = CGF.EmitScalarExpr(Device.getPointer());
+ DeviceID =
+ CGF.Builder.CreateIntCast(DevVal, CGF.Int64Ty, /*isSigned=*/true);
} else {
DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
}
@@ -12135,7 +12149,7 @@ void CGOpenMPSIMDRuntime::emitTargetOutlinedFunction(
void CGOpenMPSIMDRuntime::emitTargetCall(
CodeGenFunction &CGF, const OMPExecutableDirective &D,
llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID, const Expr *IfCond,
- const Expr *Device,
+ llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device,
llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
const OMPLoopDirective &D)>
SizeEmitter) {
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index b9df5ee11473..99b201bcf6fe 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -20,6 +20,7 @@
#include "clang/Basic/OpenMPKinds.h"
#include "clang/Basic/SourceLocation.h"
#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/PointerIntPair.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/StringMap.h"
#include "llvm/ADT/StringSet.h"
@@ -1504,16 +1505,16 @@ class CGOpenMPRuntime {
/// \param IfCond Expression evaluated in if clause associated with the target
/// directive, or null if no if clause is used.
/// \param Device Expression evaluated in device clause associated with the
- /// target directive, or null if no device clause is used.
+ /// target directive, or null if no device clause is used and device modifier.
/// \param SizeEmitter Callback to emit number of iterations for loop-based
/// directives.
- virtual void
- emitTargetCall(CodeGenFunction &CGF, const OMPExecutableDirective &D,
- llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID,
- const Expr *IfCond, const Expr *Device,
- llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
- const OMPLoopDirective &D)>
- SizeEmitter);
+ virtual void emitTargetCall(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D,
+ llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID, const Expr *IfCond,
+ llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device,
+ llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
+ const OMPLoopDirective &D)>
+ SizeEmitter);
/// Emit the target regions enclosed in \a GD function definition or
/// the function itself in case it is a valid device function. Returns true if
@@ -2275,14 +2276,14 @@ class CGOpenMPSIMDRuntime final : public CGOpenMPRuntime {
/// \param IfCond Expression evaluated in if clause associated with the target
/// directive, or null if no if clause is used.
/// \param Device Expression evaluated in device clause associated with the
- /// target directive, or null if no device clause is used.
- void
- emitTargetCall(CodeGenFunction &CGF, const OMPExecutableDirective &D,
- llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID,
- const Expr *IfCond, const Expr *Device,
- llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
- const OMPLoopDirective &D)>
- SizeEmitter) override;
+ /// target directive, or null if no device clause is used and device modifier.
+ void emitTargetCall(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D,
+ llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID, const Expr *IfCond,
+ llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device,
+ llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
+ const OMPLoopDirective &D)>
+ SizeEmitter) override;
/// Emit the target regions enclosed in \a GD function definition or
/// the function itself in case it is a valid device function. Returns true if
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 2ffe79957ee8..37bdb53bb1dc 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -4724,12 +4724,10 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
}
// Check if we have any device clause associated with the directive.
- const Expr *Device = nullptr;
- if (auto *C = S.getSingleClause<OMPDeviceClause>()) {
- if (C->getModifier() == OMPC_DEVICE_unknown ||
- C->getModifier() == OMPC_DEVICE_device_num)
- Device = C->getDevice();
- }
+ llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device(
+ nullptr, OMPC_DEVICE_unknown);
+ if (auto *C = S.getSingleClause<OMPDeviceClause>())
+ Device.setPointerAndInt(C->getDevice(), C->getModifier());
// Check if we have an if clause whose conditional always evaluates to false
// or if we do not have any targets specified. If so the target region is not
diff --git a/clang/test/OpenMP/target_device_codegen.cpp b/clang/test/OpenMP/target_device_codegen.cpp
new file mode 100644
index 000000000000..8117540d3939
--- /dev/null
+++ b/clang/test/OpenMP/target_device_codegen.cpp
@@ -0,0 +1,50 @@
+// Test host codegen.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+void foo(int n) {
+
+ // CHECK: [[N:%.+]] = load i32, i32* [[N_ADDR:%.+]],
+ // CHECK: store i32 [[N]], i32* [[DEVICE_CAP:%.+]],
+ // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
+ // CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
+ // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i64* null, i64* null)
+ // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+ // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+ // CHECK: [[FAIL]]
+ // CHECK: call void [[HVT0:@.+]]()
+ // CHECK-NEXT: br label %[[END]]
+ // CHECK: [[END]]
+ #pragma omp target device(n)
+ ;
+ // CHECK: [[N:%.+]] = load i32, i32* [[N_ADDR]],
+ // CHECK: store i32 [[N]], i32* [[DEVICE_CAP:%.+]],
+ // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
+ // CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
+ // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i64* null, i64* null)
+ // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+ // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+ // CHECK: [[FAIL]]
+ // CHECK: call void [[HVT0:@.+]]()
+ // CHECK-NEXT: br label %[[END]]
+ // CHECK: [[END]]
+ #pragma omp target device(device_num: n)
+ ;
+ // CHECK-NOT: call i32 @__tgt_target(
+ // CHECK: call void @__omp_offloading_{{.+}}_l46()
+ // CHECK-NOT: call i32 @__tgt_target(
+ #pragma omp target device(ancestor: n)
+ ;
+}
+
+#endif
More information about the cfe-commits
mailing list