forked from OSchip/llvm-project
[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.
This commit is contained in:
Alexey Bataev
parent
99336bf95a
commit
f3c857fae2
|
|
@ -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) {
|
||||
|
|
|
|||
|
|
@ -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 @@ public:
|
|||
/// \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 @@ public:
|
|||
/// \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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
Loading…
Reference in New Issue