josemonsalve2 updated this revision to Diff 359969.
josemonsalve2 added a comment.
Adding test file
`/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp`
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D106298/new/
https://reviews.llvm.org/D106298
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
Index: clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
@@ -0,0 +1,191 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK1
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK1
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK1
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK1
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK2
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK2
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK2
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK2
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK3
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK3
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK3
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK3
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK4
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK4
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK4
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK4
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK5
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK5
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK5
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK5
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK6
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK6
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK6
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK6
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK7
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK7
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK7
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=CHECK7
+
+
+// expected-no-diagnostics
+
+
+#ifndef HEADER
+#define HEADER
+
+// COM: Checking for default target and thread limit: CHECK1
+
+void default_val_num_teams() {
+ #pragma omp target
+ { int a_var; }
+
+ #pragma omp target simd
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+// COM: Check for num_teams(22) all the same: CHECK2
+void foo1() {
+ #pragma omp target teams num_teams(22)
+ { int a_var; }
+}
+
+void foo2() {
+ #pragma omp target teams distribute num_teams(22)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+void foo3() {
+ #pragma omp target teams distribute parallel for num_teams(22)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+// COM: Check for num_teams differently: CHECK3
+void bar1() {
+ #pragma omp target teams num_teams(22)
+ { int a_var; }
+}
+
+void bar2() {
+ #pragma omp target teams distribute num_teams(33)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+void bar3() {
+ #pragma omp target teams distribute parallel for num_teams(44)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+// COM: Check for const int expression: CHECK4
+void const_int() {
+ const int NT = 22;
+ #pragma omp target teams num_teams(NT)
+ { int a_var; }
+}
+
+// COM: Checking for num threads: CHECK5
+void thread_limit() {
+ #pragma omp target teams thread_limit(22)
+ { int a_var; }
+}
+
+// COM: Checking for num threads and thread limit: CHECK6
+void num_threads() {
+ #pragma omp target teams distribute parallel for thread_limit(22) num_threads(11)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+// COM: Checking for thread_limit and num_teams: CHECK6
+void threads_and_teams() {
+ #pragma omp target teams distribute parallel for thread_limit(22) num_teams(33)
+ for (int i = 0; i < 22; i++)
+ int a_var;
+}
+
+#endif
+
+
+// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}default_val_num_teams{{[^(]*}}
+// CHECK1-SAME: () #[[ATTR_OUTLINED_DEF_TEAMS:[0-9]+]] {
+// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}default_val_num_teams{{[^(]*}}
+// CHECK1-SAME: () #[[ATTR_OUTLINED_DEF_SIMD:[0-9]+]] {
+
+// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}foo1{{[^(]*}}
+// CHECK2-SAME: () #[[ATTR_OUTLINED_CHECK2:[0-9]+]] {
+// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}foo2{{[^(]*}}
+// CHECK2-SAME: () #[[ATTR_OUTLINED_CHECK2:[0-9]+]] {
+// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}foo3{{[^(]*}}
+// CHECK2-SAME: () #[[ATTR_OUTLINED_CHECK2:[0-9]+]] {
+
+// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}bar1{{[^(]*}}
+// CHECK3-SAME: () #[[ATTR_OUTLINED_CHECK3_1:[0-9]+]] {
+// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}bar2{{[^(]*}}
+// CHECK3-SAME: () #[[ATTR_OUTLINED_CHECK3_2:[0-9]+]] {
+// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}bar3{{[^(]*}}
+// CHECK3-SAME: () #[[ATTR_OUTLINED_CHECK3_3:[0-9]+]] {
+
+// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}const_int{{[^(]*}}
+// CHECK4-SAME: () #[[ATTR_OUTLINED_CHECK4:[0-9]+]] {
+
+// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}thread_limit{{[^(]*}}
+// CHECK5-SAME: () #[[ATTR_OUTLINED_CHECK5:[0-9]+]] {
+
+// CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}num_threads{{[^(]*}}
+// CHECK6-SAME: () #[[ATTR_OUTLINED_CHECK6:[0-9]+]] {
+
+// CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+.*}}threads_and_teams{{[^(]*}}
+// CHECK7-SAME: () #[[ATTR_OUTLINED_CHECK7:[0-9]+]] {
+
+
+
+// CHECK1: attributes #[[ATTR_OUTLINED_DEF_TEAMS]] = {{{.*"omp_target_num_teams"="1".*}}}
+// CHECK1: attributes #[[ATTR_OUTLINED_DEF_SIMD]] = {{{.*"omp_target_num_teams"="1".*"omp_target_thread_limit"="1".*}}}
+
+// CHECK2: attributes #[[ATTR_OUTLINED_CHECK2]] = {{{.*"omp_target_num_teams"="22".*}}}
+
+// CHECK3: attributes #[[ATTR_OUTLINED_CHECK3_1]] = {{{.*"omp_target_num_teams"="22".*}}}
+// CHECK3: attributes #[[ATTR_OUTLINED_CHECK3_2]] = {{{.*"omp_target_num_teams"="33".*}}}
+// CHECK3: attributes #[[ATTR_OUTLINED_CHECK3_3]] = {{{.*"omp_target_num_teams"="44".*}}}
+
+// CHECK4: attributes #[[ATTR_OUTLINED_CHECK4]] = {{{.*"omp_target_num_teams"="22".*}}}
+
+// CHECK5: attributes #[[ATTR_OUTLINED_CHECK5]] = {{{.*"omp_target_thread_limit"="22".*}}}
+
+// CHECK6: attributes #[[ATTR_OUTLINED_CHECK6]] = {{{.*"omp_target_thread_limit"="11".*}}}
+
+// CHECK7: attributes #[[ATTR_OUTLINED_CHECK7]] = {{{.*"omp_target_num_teams"="33".*"omp_target_thread_limit"="22".*}}}
+
+
+// void foo3() {
+// int a[N];
+// #pragma omp target teams distribute num_teams(22) thread_limit(20)
+// for (int i = 0; i < N; ++i) {
+// a[i]=5*i;
+// }
+// }
\ No newline at end of file
Index: clang/lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.h
+++ clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -340,6 +340,35 @@
llvm::Value *emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc,
unsigned Flags = 0);
+ /// Emit the number of teams for a target directive. Inspect the num_teams
+ /// clause associated with a teams construct combined or closely nested
+ /// with the target directive.
+ ///
+ /// Emit a team of size one for directives such as 'target parallel' that
+ /// have no associated teams construct.
+ ///
+ /// Otherwise, return nullptr.
+ const Expr *getNumTeamsExprForTargetDirective(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D,
+ int32_t &DefaultVal);
+ llvm::Value *emitNumTeamsForTargetDirective(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D);
+ /// Emit the number of threads for a target directive. Inspect the
+ /// thread_limit clause associated with a teams construct combined or closely
+ /// nested with the target directive.
+ ///
+ /// Emit the num_threads clause for directives such as 'target parallel' that
+ /// have no associated teams construct.
+ ///
+ /// Otherwise, return nullptr.
+ const Expr *
+ getNumThreadsExprForTargetDirective(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D,
+ int32_t &DefaultVal);
+ llvm::Value *
+ emitNumThreadsForTargetDirective(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D);
+
/// Returns pointer to ident_t type.
llvm::Type *getIdentTyPointerTy();
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6551,6 +6551,20 @@
OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID,
OffloadEntriesInfoManagerTy::OMPTargetRegionEntryTargetRegion);
+
+ // Add NumTeams and ThreadLimit attributes to the outlined GPU function
+ int32_t DefaultValTeams = -1;
+ getNumTeamsExprForTargetDirective(CGF, D, DefaultValTeams);
+ if (DefaultValTeams > 0) {
+ OutlinedFn->addFnAttr("omp_target_num_teams",
+ std::to_string(DefaultValTeams));
+ }
+ int32_t DefaultValThreads = -1;
+ getNumThreadsExprForTargetDirective(CGF, D, DefaultValThreads);
+ if (DefaultValThreads > 0) {
+ OutlinedFn->addFnAttr("omp_target_thread_limit",
+ std::to_string(DefaultValThreads));
+ }
}
/// Checks if the expression is constant or does not have non-trivial function
@@ -6605,24 +6619,13 @@
return Child;
}
-/// Emit the number of teams for a target directive. Inspect the num_teams
-/// clause associated with a teams construct combined or closely nested
-/// with the target directive.
-///
-/// Emit a team of size one for directives such as 'target parallel' that
-/// have no associated teams construct.
-///
-/// Otherwise, return nullptr.
-static llvm::Value *
-emitNumTeamsForTargetDirective(CodeGenFunction &CGF,
- const OMPExecutableDirective &D) {
- assert(!CGF.getLangOpts().OpenMPIsDevice &&
- "Clauses associated with the teams directive expected to be emitted "
- "only for the host!");
+const Expr *CGOpenMPRuntime::getNumTeamsExprForTargetDirective(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D,
+ int32_t &DefaultVal) {
+
OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
assert(isOpenMPTargetExecutionDirective(DirectiveKind) &&
"Expected target-based executable directive.");
- CGBuilderTy &Bld = CGF.Builder;
switch (DirectiveKind) {
case OMPD_target: {
const auto *CS = D.getInnermostCapturedStmt();
@@ -6638,19 +6641,22 @@
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
const Expr *NumTeams =
NestedDir->getSingleClause<OMPNumTeamsClause>()->getNumTeams();
- llvm::Value *NumTeamsVal =
- CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
- return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ if (NumTeams->isIntegerConstantExpr(CGF.getContext()))
+ if (auto Constant =
+ NumTeams->getIntegerConstantExpr(CGF.getContext()))
+ DefaultVal = Constant->getExtValue();
+ return NumTeams;
}
- return Bld.getInt32(0);
+ DefaultVal = 0;
+ return nullptr;
}
if (isOpenMPParallelDirective(NestedDir->getDirectiveKind()) ||
- isOpenMPSimdDirective(NestedDir->getDirectiveKind()))
- return Bld.getInt32(1);
- return Bld.getInt32(0);
+ isOpenMPSimdDirective(NestedDir->getDirectiveKind())) {
+ DefaultVal = 1;
+ return nullptr;
+ }
}
+ DefaultVal = 1;
return nullptr;
}
case OMPD_target_teams:
@@ -6662,19 +6668,20 @@
CodeGenFunction::RunCleanupsScope NumTeamsScope(CGF);
const Expr *NumTeams =
D.getSingleClause<OMPNumTeamsClause>()->getNumTeams();
- llvm::Value *NumTeamsVal =
- CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
- return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ if (NumTeams->isIntegerConstantExpr(CGF.getContext()))
+ if (auto Constant = NumTeams->getIntegerConstantExpr(CGF.getContext()))
+ DefaultVal = Constant->getExtValue();
+ return NumTeams;
}
- return Bld.getInt32(0);
+ DefaultVal = 0;
+ return nullptr;
}
case OMPD_target_parallel:
case OMPD_target_parallel_for:
case OMPD_target_parallel_for_simd:
case OMPD_target_simd:
- return Bld.getInt32(1);
+ DefaultVal = 1;
+ return nullptr;
case OMPD_parallel:
case OMPD_for:
case OMPD_parallel_for:
@@ -6740,6 +6747,23 @@
llvm_unreachable("Unexpected directive kind.");
}
+llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D) {
+ assert(!CGF.getLangOpts().OpenMPIsDevice &&
+ "Clauses associated with the teams directive expected to be emitted "
+ "only for the host!");
+ CGBuilderTy &Bld = CGF.Builder;
+ int32_t DefaultNT = 0;
+ const Expr *NumTeams = getNumTeamsExprForTargetDirective(CGF, D, DefaultNT);
+ if (NumTeams != nullptr) {
+ llvm::Value *NumTeamsVal = CGF.EmitScalarExpr(NumTeams,
+ /*IgnoreResultAssign*/ true);
+ return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
+ /*isSigned=*/true);
+ }
+ return Bld.getInt32(DefaultNT);
+}
+
static llvm::Value *getNumThreads(CodeGenFunction &CGF, const CapturedStmt *CS,
llvm::Value *DefaultThreadLimitVal) {
const Stmt *Child = CGOpenMPRuntime::getSingleCompoundChild(
@@ -6832,17 +6856,130 @@
: CGF.Builder.getInt32(0);
}
-/// Emit the number of threads for a target directive. Inspect the
-/// thread_limit clause associated with a teams construct combined or closely
-/// nested with the target directive.
-///
-/// Emit the num_threads clause for directives such as 'target parallel' that
-/// have no associated teams construct.
-///
-/// Otherwise, return nullptr.
-static llvm::Value *
-emitNumThreadsForTargetDirective(CodeGenFunction &CGF,
- const OMPExecutableDirective &D) {
+const Expr *CGOpenMPRuntime::getNumThreadsExprForTargetDirective(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D,
+ int32_t &DefaultVal) {
+ OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
+ assert(isOpenMPTargetExecutionDirective(DirectiveKind) &&
+ "Expected target-based executable directive.");
+
+ switch (DirectiveKind) {
+ case OMPD_target:
+ // Teams have no clause thread_limit
+ return nullptr;
+ case OMPD_target_teams:
+ case OMPD_target_teams_distribute:
+ if (D.hasClausesOfKind<OMPThreadLimitClause>()) {
+ const auto *ThreadLimitClause = D.getSingleClause<OMPThreadLimitClause>();
+ const Expr *ThreadLimit = ThreadLimitClause->getThreadLimit();
+ if (ThreadLimit->isIntegerConstantExpr(CGF.getContext()))
+ if (auto Constant =
+ ThreadLimit->getIntegerConstantExpr(CGF.getContext()))
+ DefaultVal = Constant->getExtValue();
+ return ThreadLimit;
+ }
+ return nullptr;
+ case OMPD_target_parallel:
+ case OMPD_target_parallel_for:
+ case OMPD_target_parallel_for_simd:
+ case OMPD_target_teams_distribute_parallel_for:
+ case OMPD_target_teams_distribute_parallel_for_simd: {
+ Expr *ThreadLimit = nullptr;
+ Expr *NumThreads = nullptr;
+ if (D.hasClausesOfKind<OMPThreadLimitClause>()) {
+ const auto *ThreadLimitClause = D.getSingleClause<OMPThreadLimitClause>();
+ ThreadLimit = ThreadLimitClause->getThreadLimit();
+ if (ThreadLimit->isIntegerConstantExpr(CGF.getContext()))
+ if (auto Constant =
+ ThreadLimit->getIntegerConstantExpr(CGF.getContext()))
+ DefaultVal = Constant->getExtValue();
+ }
+ if (D.hasClausesOfKind<OMPNumThreadsClause>()) {
+ const auto *NumThreadsClause = D.getSingleClause<OMPNumThreadsClause>();
+ NumThreads = NumThreadsClause->getNumThreads();
+ if (NumThreads->isIntegerConstantExpr(CGF.getContext())) {
+ if (auto Constant =
+ NumThreads->getIntegerConstantExpr(CGF.getContext())) {
+ if (Constant->getExtValue() < DefaultVal) {
+ DefaultVal = Constant->getExtValue();
+ ThreadLimit = NumThreads;
+ }
+ }
+ }
+ }
+ return ThreadLimit;
+ }
+ case OMPD_target_teams_distribute_simd:
+ case OMPD_target_simd:
+ DefaultVal = 1;
+ return nullptr;
+ case OMPD_parallel:
+ case OMPD_for:
+ case OMPD_parallel_for:
+ case OMPD_parallel_master:
+ case OMPD_parallel_sections:
+ case OMPD_for_simd:
+ case OMPD_parallel_for_simd:
+ case OMPD_cancel:
+ case OMPD_cancellation_point:
+ case OMPD_ordered:
+ case OMPD_threadprivate:
+ case OMPD_allocate:
+ case OMPD_task:
+ case OMPD_simd:
+ case OMPD_tile:
+ case OMPD_unroll:
+ case OMPD_sections:
+ case OMPD_section:
+ case OMPD_single:
+ case OMPD_master:
+ case OMPD_critical:
+ case OMPD_taskyield:
+ case OMPD_barrier:
+ case OMPD_taskwait:
+ case OMPD_taskgroup:
+ case OMPD_atomic:
+ case OMPD_flush:
+ case OMPD_depobj:
+ case OMPD_scan:
+ case OMPD_teams:
+ case OMPD_target_data:
+ case OMPD_target_exit_data:
+ case OMPD_target_enter_data:
+ case OMPD_distribute:
+ case OMPD_distribute_simd:
+ case OMPD_distribute_parallel_for:
+ case OMPD_distribute_parallel_for_simd:
+ case OMPD_teams_distribute:
+ case OMPD_teams_distribute_simd:
+ case OMPD_teams_distribute_parallel_for:
+ case OMPD_teams_distribute_parallel_for_simd:
+ case OMPD_target_update:
+ case OMPD_declare_simd:
+ case OMPD_declare_variant:
+ case OMPD_begin_declare_variant:
+ case OMPD_end_declare_variant:
+ case OMPD_declare_target:
+ case OMPD_end_declare_target:
+ case OMPD_declare_reduction:
+ case OMPD_declare_mapper:
+ case OMPD_taskloop:
+ case OMPD_taskloop_simd:
+ case OMPD_master_taskloop:
+ case OMPD_master_taskloop_simd:
+ case OMPD_parallel_master_taskloop:
+ case OMPD_parallel_master_taskloop_simd:
+ case OMPD_requires:
+ case OMPD_unknown:
+ break;
+ default:
+ break;
+ }
+ llvm_unreachable("Unsupported directive kind.");
+}
+
+llvm::Value *CGOpenMPRuntime::emitNumThreadsForTargetDirective(
+ CodeGenFunction &CGF, const OMPExecutableDirective &D) {
assert(!CGF.getLangOpts().OpenMPIsDevice &&
"Clauses associated with the teams directive expected to be emitted "
"only for the host!");
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
