diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/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(); diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/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()->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()->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()) { + const auto *ThreadLimitClause = D.getSingleClause(); + 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()) { + const auto *ThreadLimitClause = D.getSingleClause(); + ThreadLimit = ThreadLimitClause->getThreadLimit(); + if (ThreadLimit->isIntegerConstantExpr(CGF.getContext())) + if (auto Constant = + ThreadLimit->getIntegerConstantExpr(CGF.getContext())) + DefaultVal = Constant->getExtValue(); + } + if (D.hasClausesOfKind()) { + const auto *NumThreadsClause = D.getSingleClause(); + 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!"); diff --git a/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp b/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp new file mode 100644 --- /dev/null +++ b/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