From: Alexey Bataev Date: Wed, 21 Nov 2018 21:04:34 +0000 (+0000) Subject: [OPENMP][NVPTX]Emit default locations as constant with undefined mode. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=5e239bf6b632486135b28f2784aed77765b4c1c9;p=clang [OPENMP][NVPTX]Emit default locations as constant with undefined mode. For the NVPTX target default locations should be emitted as constants + additional info must be emitted in the reserved_2 field of the ident_t structure. The 1st bit controls the execution mode and the 2nd bit controls use of the lightweight runtime. The combination of the bits for Non-SPMD mode + lightweight runtime represents special undefined mode, used outside of the target regions for orphaned directives or functions. Should allow and additional optimization inside of the target regions. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@347425 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp index cf2c6cf6f2..cade093941 100644 --- a/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/lib/CodeGen/CGOpenMPRuntime.cpp @@ -1467,7 +1467,9 @@ createConstantGlobalStructAndAddToParent(CodeGenModule &CGM, QualType Ty, Address CGOpenMPRuntime::getOrCreateDefaultLocation(unsigned Flags) { CharUnits Align = CGM.getContext().getTypeAlignInChars(IdentQTy); - llvm::Value *Entry = OpenMPDefaultLocMap.lookup(Flags); + unsigned Reserved2Flags = getDefaultLocationReserved2Flags(); + FlagsTy FlagsKey(Flags, Reserved2Flags); + llvm::Value *Entry = OpenMPDefaultLocMap.lookup(FlagsKey); if (!Entry) { if (!DefaultOpenMPPSource) { // Initialize default location for psource field of ident_t structure of @@ -1480,18 +1482,18 @@ Address CGOpenMPRuntime::getOrCreateDefaultLocation(unsigned Flags) { llvm::ConstantExpr::getBitCast(DefaultOpenMPPSource, CGM.Int8PtrTy); } - llvm::Constant *Data[] = {llvm::ConstantInt::getNullValue(CGM.Int32Ty), - llvm::ConstantInt::get(CGM.Int32Ty, Flags), - llvm::ConstantInt::getNullValue(CGM.Int32Ty), - llvm::ConstantInt::getNullValue(CGM.Int32Ty), - DefaultOpenMPPSource}; + llvm::Constant *Data[] = { + llvm::ConstantInt::getNullValue(CGM.Int32Ty), + llvm::ConstantInt::get(CGM.Int32Ty, Flags), + llvm::ConstantInt::get(CGM.Int32Ty, Reserved2Flags), + llvm::ConstantInt::getNullValue(CGM.Int32Ty), DefaultOpenMPPSource}; llvm::GlobalValue *DefaultOpenMPLocation = - createGlobalStruct(CGM, IdentQTy, /*IsConstant=*/false, Data, "", + createGlobalStruct(CGM, IdentQTy, isDefaultLocationConstant(), Data, "", llvm::GlobalValue::PrivateLinkage); DefaultOpenMPLocation->setUnnamedAddr( llvm::GlobalValue::UnnamedAddr::Global); - OpenMPDefaultLocMap[Flags] = Entry = DefaultOpenMPLocation; + OpenMPDefaultLocMap[FlagsKey] = Entry = DefaultOpenMPLocation; } return Address(Entry, Align); } diff --git a/lib/CodeGen/CGOpenMPRuntime.h b/lib/CodeGen/CGOpenMPRuntime.h index 1a27db1574..489ecdf15e 100644 --- a/lib/CodeGen/CGOpenMPRuntime.h +++ b/lib/CodeGen/CGOpenMPRuntime.h @@ -282,12 +282,21 @@ protected: bool AtCurrentPoint = false); void clearLocThreadIdInsertPt(CodeGenFunction &CGF); + /// Check if the default location must be constant. + /// Default is false to support OMPT/OMPD. + virtual bool isDefaultLocationConstant() const { return false; } + + /// Returns additional flags that can be stored in reserved_2 field of the + /// default location. + virtual unsigned getDefaultLocationReserved2Flags() const { return 0; } + private: /// Default const ident_t object used for initialization of all other /// ident_t objects. llvm::Constant *DefaultOpenMPPSource = nullptr; + using FlagsTy = std::pair; /// Map of flags and corresponding default locations. - typedef llvm::DenseMap OpenMPDefaultLocMapTy; + using OpenMPDefaultLocMapTy = llvm::DenseMap; OpenMPDefaultLocMapTy OpenMPDefaultLocMap; Address getOrCreateDefaultLocation(unsigned Flags); diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp index 74b24eacce..60fbf66811 100644 --- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp +++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp @@ -1902,6 +1902,26 @@ void CGOpenMPRuntimeNVPTX::emitTargetOutlinedFunction( setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode); } +namespace { +LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE(); +/// Enum for accesseing the reserved_2 field of the ident_t struct. +enum ModeFlagsTy : unsigned { + /// Bit set to 1 when in SPMD mode. + KMP_IDENT_SPMD_MODE = 0x01, + /// Bit set to 1 when a simplified runtime is used. + KMP_IDENT_SIMPLE_RT_MODE = 0x02, + LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/KMP_IDENT_SIMPLE_RT_MODE) +}; + +/// Special mode Undefined. Is the combination of Non-SPMD mode + SimpleRuntime. +static const ModeFlagsTy UndefinedMode = + (~KMP_IDENT_SPMD_MODE) & KMP_IDENT_SIMPLE_RT_MODE; +} // anonymous namespace + +unsigned CGOpenMPRuntimeNVPTX::getDefaultLocationReserved2Flags() const { + return UndefinedMode; +} + CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM) : CGOpenMPRuntime(CGM, "_", "$") { if (!CGM.getLangOpts().OpenMPIsDevice) diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h index 2acab07358..b03ff782d7 100644 --- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h +++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h @@ -180,6 +180,16 @@ protected: return "__omp_outlined__"; } + /// Check if the default location must be constant. + /// Constant for NVPTX for better optimization. + bool isDefaultLocationConstant() const override { return true; } + + /// Returns additional flags that can be stored in reserved_2 field of the + /// default location. + /// For NVPTX target contains data about SPMD/Non-SPMD execution mode + + /// Full/Lightweight runtime mode. Used for better optimization. + unsigned getDefaultLocationReserved2Flags() const override; + public: explicit CGOpenMPRuntimeNVPTX(CodeGenModule &CGM); void clear() override; diff --git a/test/OpenMP/nvptx_SPMD_codegen.cpp b/test/OpenMP/nvptx_SPMD_codegen.cpp index cf68ee7dd8..97481e8d67 100644 --- a/test/OpenMP/nvptx_SPMD_codegen.cpp +++ b/test/OpenMP/nvptx_SPMD_codegen.cpp @@ -8,6 +8,11 @@ #ifndef HEADER #define HEADER +// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 +// CHECK-DAG: private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 2, i32 0, i8* getelementptr inbounds +// CHECK-DAG: private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 2, i32 0, i8* getelementptr inbounds +// CHECK-DAG: private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 2, i32 0, i8* getelementptr inbounds +// CHECK-DAG: private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 2, i32 0, i8* getelementptr inbounds // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 void foo() { diff --git a/test/OpenMP/nvptx_target_printf_codegen.c b/test/OpenMP/nvptx_target_printf_codegen.c index e7bfb874f4..098c8e165f 100644 --- a/test/OpenMP/nvptx_target_printf_codegen.c +++ b/test/OpenMP/nvptx_target_printf_codegen.c @@ -6,8 +6,10 @@ // expected-no-diagnostics extern int printf(const char *, ...); +// CHECK-DAG: private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 2, i32 0, i8* getelementptr inbounds + // Check a simple call to printf end-to-end. -// CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double } +// CHECK-DAG: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double } int CheckSimple() { // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+CheckSimple.+]]_worker() #pragma omp target