shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[OpenMP] Introduce the initial support for OpenMP kernel language (#66844)

Browse Source 
This patch starts the support for OpenMP kernel language, basically to
write
OpenMP target region in SIMT style, similar to kernel languages such as
CUDA.
What included in this first patch is the `ompx_bare` clause for `target
teams`
directive. When `ompx_bare` exists, globalization is disabled such that
local
variables will not be globalized. The runtime init/deinit function calls
will
not be emitted. That being said, almost all OpenMP executable directives
are
not supported in the region, such as parallel, task. This patch doesn't
include
the Sema checks for that, so the use of them is UB. Simple directives,
such as
atomic, can be used. We provide a set of APIs (for C, they are prefix
with
`ompx_`; for C++, they are in `ompx` namespace) to get thread id, block
id, etc.
For more details, you can refer to
https://tianshilei.me/wp-content/uploads/llvm-hpc-2023.pdf.
This commit is contained in:
Shilei Tian 2023-09-29 13:11:09 -04:00 committed by GitHub
parent 2fc90afdac
commit e997dca333
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
20 changed files with 1145 additions and 657 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
clang/include/clang/AST/OpenMPClause.h
View File

@ -9220,6 +9220,27 @@ private:
}
};
/// This represents 'ompx_bare' clause in the '#pragma omp target teams ...'
/// directive.
///
/// \code
/// #pragma omp target teams ompx_bare
/// \endcode
/// In this example directive '#pragma omp target teams' has a 'ompx_bare'
/// clause.
class OMPXBareClause : public OMPNoChildClause<llvm::omp::OMPC_ompx_bare> {
public:
/// Build 'ompx_bare' clause.
///
/// \param StartLoc Starting location of the clause.
/// \param EndLoc Ending location of the clause.
OMPXBareClause(SourceLocation StartLoc, SourceLocation EndLoc)
: OMPNoChildClause(StartLoc, EndLoc) {}
/// Build an empty clause.
OMPXBareClause() = default;
};
} // namespace clang
#endif // LLVM_CLANG_AST_OPENMPCLAUSE_H

5
clang/include/clang/AST/RecursiveASTVisitor.h
View File

@ -3890,6 +3890,11 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXAttributeClause(
return true;
}
template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPXBareClause(OMPXBareClause *C) {
return true;
}
// FIXME: look at the following tricky-seeming exprs to see if we
// need to recurse on anything. These are ones that have methods
// returning decls or qualtypes or nestednamespecifier -- though I'm

4
clang/include/clang/Basic/DiagnosticParseKinds.td
View File

@ -1360,6 +1360,8 @@ def warn_clause_expected_string : Warning<
"expected string literal in 'clause %0' - ignoring">, InGroup<IgnoredPragmas>;
def err_omp_unexpected_clause : Error<
"unexpected OpenMP clause '%0' in directive '#pragma omp %1'">;
def err_omp_unexpected_clause_extension_only : Error<
"OpenMP clause '%0' is only available as extension, use '-fopenmp-extensions'">;
def err_omp_immediate_directive : Error<
"'#pragma omp %0' %select{|with '%2' clause }1cannot be an immediate substatement">;
def err_omp_expected_identifier_for_critical : Error<
@ -1452,6 +1454,8 @@ def warn_unknown_declare_variant_isa_trait
"spelling or consider restricting the context selector with the "
"'arch' selector further">,
InGroup<SourceUsesOpenMP>;
def note_ompx_bare_clause : Note<
"OpenMP extension clause '%0' only allowed with '#pragma omp %1'">;
def note_omp_declare_variant_ctx_options
: Note<"context %select{set|selector|property}0 options are: %1">;
def warn_omp_declare_variant_expected

4
clang/include/clang/Sema/Sema.h
View File

@ -12448,6 +12448,10 @@ public:
SourceLocation LParenLoc,
SourceLocation EndLoc);
/// Called on a well-formed 'ompx_bare' clause.
OMPClause *ActOnOpenMPXBareClause(SourceLocation StartLoc,
SourceLocation EndLoc);
/// The kind of conversion being performed.
enum CheckedConversionKind {
/// An implicit conversion.

5
clang/lib/AST/OpenMPClause.cpp
View File

@ -170,6 +170,7 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
case OMPC_affinity:
case OMPC_when:
case OMPC_bind:
case OMPC_ompx_bare:
break;
default:
break;
@ -2546,6 +2547,10 @@ void OMPClausePrinter::VisitOMPXAttributeClause(OMPXAttributeClause *Node) {
OS << ")";
}
void OMPClausePrinter::VisitOMPXBareClause(OMPXBareClause *Node) {
OS << "ompx_bare";
}
void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
VariantMatchInfo &VMI) const {
for (const OMPTraitSet &Set : Sets) {

1
clang/lib/AST/StmtProfile.cpp
View File

@ -930,6 +930,7 @@ void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
}
void OMPClauseProfiler::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
}
void OMPClauseProfiler::VisitOMPXBareClause(const OMPXBareClause *C) {}
} // namespace
void

59
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
View File

@ -551,10 +551,9 @@ CGOpenMPRuntimeGPU::getExecutionMode() const {
return CurrentExecutionMode;
}
static CGOpenMPRuntimeGPU::DataSharingMode
getDataSharingMode(CodeGenModule &CGM) {
return CGM.getLangOpts().OpenMPCUDAMode ? CGOpenMPRuntimeGPU::CUDA
: CGOpenMPRuntimeGPU::Generic;
CGOpenMPRuntimeGPU::DataSharingMode
CGOpenMPRuntimeGPU::getDataSharingMode() const {
return CurrentDataSharingMode;
}
/// Check for inner (nested) SPMD construct, if any
@ -752,6 +751,9 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
EntryFunctionState EST;
WrapperFunctionsMap.clear();
[[maybe_unused]] bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
assert(!IsBareKernel && "bare kernel should not be at generic mode");
// Emit target region as a standalone region.
class NVPTXPrePostActionTy : public PrePostActionTy {
CGOpenMPRuntimeGPU::EntryFunctionState &EST;
@ -760,15 +762,13 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST)
: EST(EST) {}
void Enter(CodeGenFunction &CGF) override {
auto &RT =
static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
RT.emitKernelInit(CGF, EST, /* IsSPMD */ false);
// Skip target region initialization.
RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
}
void Exit(CodeGenFunction &CGF) override {
auto &RT =
static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
RT.clearLocThreadIdInsertPt(CGF);
RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);
}
@ -807,25 +807,39 @@ void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);
EntryFunctionState EST;
bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
// Emit target region as a standalone region.
class NVPTXPrePostActionTy : public PrePostActionTy {
CGOpenMPRuntimeGPU &RT;
CGOpenMPRuntimeGPU::EntryFunctionState &EST;
bool IsBareKernel;
DataSharingMode Mode;
public:
NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
CGOpenMPRuntimeGPU::EntryFunctionState &EST)
: RT(RT), EST(EST) {}
CGOpenMPRuntimeGPU::EntryFunctionState &EST,
bool IsBareKernel)
: RT(RT), EST(EST), IsBareKernel(IsBareKernel),
Mode(RT.CurrentDataSharingMode) {}
void Enter(CodeGenFunction &CGF) override {
if (IsBareKernel) {
RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;
return;
}
RT.emitKernelInit(CGF, EST, /* IsSPMD */ true);
// Skip target region initialization.
RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
}
void Exit(CodeGenFunction &CGF) override {
if (IsBareKernel) {
RT.CurrentDataSharingMode = Mode;
return;
}
RT.clearLocThreadIdInsertPt(CGF);
RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
}
} Action(*this, EST);
} Action(*this, EST, IsBareKernel);
CodeGen.setAction(Action);
IsInTTDRegion = true;
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
@ -843,7 +857,8 @@ void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
assert(!ParentName.empty() && "Invalid target region parent name!");
bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);
if (Mode)
bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
if (Mode || IsBareKernel)
emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
CodeGen);
else
@ -867,6 +882,9 @@ CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
if (CGM.getLangOpts().NoGPULib || CGM.getLangOpts().OMPHostIRFile.empty())
return;
if (CGM.getLangOpts().OpenMPCUDAMode)
CurrentDataSharingMode = CGOpenMPRuntimeGPU::DS_CUDA;
OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTargetDebug,
"__omp_rtl_debug_kind");
OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTeamSubscription,
@ -1030,7 +1048,7 @@ llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
SourceLocation Loc,
bool WithSPMDCheck) {
if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic &&
getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
return;
@ -1142,7 +1160,7 @@ void CGOpenMPRuntimeGPU::getKmpcFreeShared(
void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF,
bool WithSPMDCheck) {
if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic &&
getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
return;
@ -1178,11 +1196,18 @@ void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
if (!CGF.HaveInsertPoint())
return;
bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
/*Name=*/".zero.addr");
CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);
llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).getPointer());
// We don't emit any thread id function call in bare kernel, but because the
// outlined function has a pointer argument, we emit a nullptr here.
if (IsBareKernel)
OutlinedFnArgs.push_back(llvm::ConstantPointerNull::get(CGM.VoidPtrTy));
else
OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).getPointer());
OutlinedFnArgs.push_back(ZeroAddr.getPointer());
OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);
@ -3273,7 +3298,7 @@ llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
const Decl *D) {
if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
return;
assert(D && "Expected function or captured|block decl.");
@ -3382,7 +3407,7 @@ Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
VarTy, Align);
}
if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
return Address::invalid();
VD = VD->getCanonicalDecl();

29
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
View File

@ -32,6 +32,18 @@ public:
/// Unknown execution mode (orphaned directive).
EM_Unknown,
};
/// Target codegen is specialized based on two data-sharing modes: CUDA, in
/// which the local variables are actually global threadlocal, and Generic, in
/// which the local variables are placed in global memory if they may escape
/// their declaration context.
enum DataSharingMode {
/// CUDA data sharing mode.
DS_CUDA,
/// Generic data-sharing mode.
DS_Generic,
};
private:
/// Parallel outlined function work for workers to execute.
llvm::SmallVector<llvm::Function *, 16> Work;
@ -42,6 +54,8 @@ private:
ExecutionMode getExecutionMode() const;
DataSharingMode getDataSharingMode() const;
/// Get barrier to synchronize all threads in a block.
void syncCTAThreads(CodeGenFunction &CGF);
@ -297,17 +311,6 @@ public:
Address getAddressOfLocalVariable(CodeGenFunction &CGF,
const VarDecl *VD) override;
/// Target codegen is specialized based on two data-sharing modes: CUDA, in
/// which the local variables are actually global threadlocal, and Generic, in
/// which the local variables are placed in global memory if they may escape
/// their declaration context.
enum DataSharingMode {
/// CUDA data sharing mode.
CUDA,
/// Generic data-sharing mode.
Generic,
};
/// Cleans up references to the objects in finished function.
///
void functionFinished(CodeGenFunction &CGF) override;
@ -343,6 +346,10 @@ private:
/// to emit optimized code.
ExecutionMode CurrentExecutionMode = EM_Unknown;
/// Track the data sharing mode when codegening directives within a target
/// region.
DataSharingMode CurrentDataSharingMode = DataSharingMode::DS_Generic;
/// true if currently emitting code for target/teams/distribute region, false
/// - otherwise.
bool IsInTTDRegion = false;

11
clang/lib/Parse/ParseOpenMP.cpp
View File

@ -3416,6 +3416,17 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
case OMPC_ompx_attribute:
Clause = ParseOpenMPOMPXAttributesClause(WrongDirective);
break;
case OMPC_ompx_bare:
if (WrongDirective)
Diag(Tok, diag::note_ompx_bare_clause)
<< getOpenMPClauseName(CKind) << "target teams";
if (!ErrorFound && !getLangOpts().OpenMPExtensions) {
Diag(Tok, diag::err_omp_unexpected_clause_extension_only)
<< getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind);
ErrorFound = true;
}
Clause = ParseOpenMPClause(CKind, WrongDirective);
break;
default:
break;
}

8
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -17553,6 +17553,9 @@ OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
case OMPC_partial:
Res = ActOnOpenMPPartialClause(nullptr, StartLoc, /*LParenLoc=*/{}, EndLoc);
break;
case OMPC_ompx_bare:
Res = ActOnOpenMPXBareClause(StartLoc, EndLoc);
break;
case OMPC_if:
case OMPC_final:
case OMPC_num_threads:
@ -24279,3 +24282,8 @@ OMPClause *Sema::ActOnOpenMPXAttributeClause(ArrayRef<const Attr *> Attrs,
SourceLocation EndLoc) {
return new (Context) OMPXAttributeClause(Attrs, StartLoc, LParenLoc, EndLoc);
}
OMPClause *Sema::ActOnOpenMPXBareClause(SourceLocation StartLoc,
SourceLocation EndLoc) {
return new (Context) OMPXBareClause(StartLoc, EndLoc);
}

14
clang/lib/Sema/TreeTransform.h
View File

@ -2391,6 +2391,15 @@ public:
EndLoc);
}
/// Build a new OpenMP 'ompx_bare' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
/// Subclasses may override this routine to provide different behavior.
OMPClause *RebuildOMPXBareClause(SourceLocation StartLoc,
SourceLocation EndLoc) {
return getSema().ActOnOpenMPXBareClause(StartLoc, EndLoc);
}
/// Build a new OpenMP 'align' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
@ -10804,6 +10813,11 @@ TreeTransform<Derived>::TransformOMPXAttributeClause(OMPXAttributeClause *C) {
NewAttrs, C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
}
template <typename Derived>
OMPClause *TreeTransform<Derived>::TransformOMPXBareClause(OMPXBareClause *C) {
return getDerived().RebuildOMPXBareClause(C->getBeginLoc(), C->getEndLoc());
}
//===----------------------------------------------------------------------===//
// Expression transformation
//===----------------------------------------------------------------------===//

5
clang/lib/Serialization/ASTReader.cpp
View File

@ -10446,6 +10446,9 @@ OMPClause *OMPClauseReader::readClause() {
case llvm::omp::OMPC_ompx_attribute:
C = new (Context) OMPXAttributeClause();
break;
case llvm::omp::OMPC_ompx_bare:
C = new (Context) OMPXBareClause();
break;
#define OMP_CLAUSE_NO_CLASS(Enum, Str) \
case llvm::omp::Enum: \
break;
@ -11547,6 +11550,8 @@ void OMPClauseReader::VisitOMPXAttributeClause(OMPXAttributeClause *C) {
C->setLocEnd(Record.readSourceLocation());
}
void OMPClauseReader::VisitOMPXBareClause(OMPXBareClause *C) {}
OMPTraitInfo *ASTRecordReader::readOMPTraitInfo() {
OMPTraitInfo &TI = getContext().getNewOMPTraitInfo();
TI.Sets.resize(readUInt32());

2
clang/lib/Serialization/ASTWriter.cpp
View File

@ -7258,6 +7258,8 @@ void OMPClauseWriter::VisitOMPXAttributeClause(OMPXAttributeClause *C) {
Record.AddSourceLocation(C->getEndLoc());
}
void OMPClauseWriter::VisitOMPXBareClause(OMPXBareClause *C) {}
void ASTRecordWriter::writeOMPTraitInfo(const OMPTraitInfo *TI) {
writeUInt32(TI->Sets.size());
for (const auto &Set : TI->Sets) {

56
clang/test/OpenMP/nvptx_target_teams_ompx_bare_codegen.cpp Normal file
View File

@ -0,0 +1,56 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
template<typename tx>
tx ftemplate(int n) {
tx a = 0;
#pragma omp target teams ompx_bare
{
a = 2;
}
return a;
}
int bar(int n){
int a = 0;
a += ftemplate<char>(n);
return a;
}
#endif
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l13
// CHECK-SAME: (i64 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8
// CHECK-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8
// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
// CHECK-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[A_ADDR]], align 1
// CHECK-NEXT: store i8 [[TMP0]], ptr [[A_CASTED]], align 1
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l13_omp_outlined(ptr null, ptr [[DOTZERO_ADDR]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]]
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l13_omp_outlined
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
// CHECK-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8
// CHECK-NEXT: store i8 2, ptr [[A_ADDR]], align 1
// CHECK-NEXT: ret void
//

21
clang/test/OpenMP/ompx_bare_messages.c Normal file
View File

@ -0,0 +1,21 @@
// RUN: %clang_cc1 -verify -fopenmp %s
// RUN: %clang_cc1 -verify -fopenmp-simd %s
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64 %s
void foo() {
}
void bar() {
#pragma omp target ompx_bare // expected-error {{unexpected OpenMP clause 'ompx_bare' in directive '#pragma omp target'}} expected-note {{OpenMP extension clause 'ompx_bare' only allowed with '#pragma omp target teams'}}
foo();
#pragma omp target teams distribute ompx_bare // expected-error {{unexpected OpenMP clause 'ompx_bare' in directive '#pragma omp target teams distribute'}} expected-note {{OpenMP extension clause 'ompx_bare' only allowed with '#pragma omp target teams'}}
for (int i = 0; i < 10; ++i) {}
#pragma omp target teams distribute parallel for ompx_bare // expected-error {{unexpected OpenMP clause 'ompx_bare' in directive '#pragma omp target teams distribute parallel for'}} expected-note {{OpenMP extension clause 'ompx_bare' only allowed with '#pragma omp target teams'}}
for (int i = 0; i < 10; ++i) {}
#pragma omp target
#pragma omp teams ompx_bare // expected-error {{unexpected OpenMP clause 'ompx_bare' in directive '#pragma omp teams'}} expected-note {{OpenMP extension clause 'ompx_bare' only allowed with '#pragma omp target teams'}}
foo();
}

4
clang/test/OpenMP/target_teams_ast_print.cpp
View File

@ -111,6 +111,10 @@ int main (int argc, char **argv) {
// CHECK-NEXT: #pragma omp target teams
a=2;
// CHECK-NEXT: a = 2;
#pragma omp target teams ompx_bare
// CHECK-NEXT: #pragma omp target teams ompx_bare
a=3;
// CHECK-NEXT: a = 3;
#pragma omp target teams default(none), private(argc,b) num_teams(f) firstprivate(argv) reduction(| : c, d) reduction(* : e) thread_limit(f+g)
// CHECK-NEXT: #pragma omp target teams default(none) private(argc,b) num_teams(f) firstprivate(argv) reduction(|: c,d) reduction(*: e) thread_limit(f + g)
foo();

1546
clang/test/OpenMP/target_teams_codegen.cpp
View File

File diff suppressed because it is too large Load Diff

1
clang/tools/libclang/CIndex.cpp
View File

@ -2735,6 +2735,7 @@ void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
}
void OMPClauseEnqueue::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
}
void OMPClauseEnqueue::VisitOMPXBareClause(const OMPXBareClause *C) {}
} // namespace

1
flang/lib/Semantics/check-omp-structure.cpp
View File

@ -2148,6 +2148,7 @@ CHECK_SIMPLE_CLAUSE(Compare, OMPC_compare)
CHECK_SIMPLE_CLAUSE(CancellationConstructType, OMPC_cancellation_construct_type)
CHECK_SIMPLE_CLAUSE(Doacross, OMPC_doacross)
CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute)
CHECK_SIMPLE_CLAUSE(OmpxBare, OMPC_ompx_bare)
CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize)
CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks)

5
llvm/include/llvm/Frontend/OpenMP/OMP.td
View File

@ -450,6 +450,10 @@ def OMPC_OMPX_Attribute : Clause<"ompx_attribute"> {
let clangClass = "OMPXAttributeClause";
}
def OMPC_OMX_Bare : Clause<"ompx_bare"> {
let clangClass = "OMPXBareClause";
}
//===----------------------------------------------------------------------===//
// Definition of OpenMP directives
//===----------------------------------------------------------------------===//
@ -1520,6 +1524,7 @@ def OMP_TargetTeams : Directive<"target teams"> {
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_ThreadLimit>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_OMX_Bare>,
];
}
def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> {