[clang] [Clang][OpenMP] Fix `if` clause not affecting `teams` construct (PR #207444)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 3 11:31:49 PDT 2026
Jan =?utf-8?q?André?= Reuter <jan at zyten.de>
Message-ID:
In-Reply-To: <llvm.org/llvm/llvm-project/pull/207444 at github.com>
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-clang
Author: Jan André Reuter (Thyre)
<details>
<summary>Changes</summary>
In https://github.com/llvm/llvm-project/issues/205345, it was discovered that the `if` clause for the `teams` construct
does not actually create a serial league when the code is built with Clang.
Looking at the emitted IR, no call to `__kmpc_push_num_threads` was added
for an `if` clause. Hence, `__kmpc_fork_teams` always created the number of
teams specified via the `num_teams` clause, or the _nteams-var_ ICV.
During `emitCommonOMPTeamsDirective`, only the `num_teams` and `thread_limit`
clauses were checked, whereas the `if` clause was missing.
To fix the missing code, check for an `if` clause mapped to `teams` or
`unknown`. If the clause is present, and true, emit `__kmpc_push_num_teams`
with the number of teams set to 1. The `thread_limit` clause is unaffected.
Otherwise, stick to the existing implementation and only call
`__kmpc_push_num_teams` if `num_teams` or `thread_limit` are set.
This only applies for OpenMP v5.2 and newer. Earlier OpenMP versions did not
support the `if` clause for the `teams` construct, hence should continue
to emit a number of teams based on attached `num_teams` and `thread_limit`
clauses, or the _nteams-var_ ICV.
---
Patch is 1.66 MiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/207444.diff
3 Files Affected:
- (modified) clang/lib/CodeGen/CGStmtOpenMP.cpp (+41-8)
- (added) clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_openmp52_codegen.cpp (+12098)
- (added) clang/test/OpenMP/teams_distribute_parallel_for_simd_if_openmp52_codegen.cpp (+12190)
``````````diff
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 88f698f38cce0..d90099d357cba 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -7366,14 +7366,47 @@ static void emitCommonOMPTeamsDirective(CodeGenFunction &CGF,
CGF, S, *CS->getCapturedDecl()->param_begin(), InnermostKind,
CodeGen);
- const auto *NT = S.getSingleClause<OMPNumTeamsClause>();
- const auto *TL = S.getSingleClause<OMPThreadLimitClause>();
- if (NT || TL) {
- const Expr *NumTeams = NT ? NT->getNumTeams().front() : nullptr;
- const Expr *ThreadLimit = TL ? TL->getThreadLimit().front() : nullptr;
-
- CGF.CGM.getOpenMPRuntime().emitNumTeamsClause(CGF, NumTeams, ThreadLimit,
- S.getBeginLoc());
+ auto &&ParallelLeague = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+ const auto *NT = S.getSingleClause<OMPNumTeamsClause>();
+ const auto *TL = S.getSingleClause<OMPThreadLimitClause>();
+ if (NT || TL) {
+ const Expr *NumTeams = NT ? NT->getNumTeams().front() : nullptr;
+ const Expr *ThreadLimit = TL ? TL->getThreadLimit().front() : nullptr;
+
+ CGF.CGM.getOpenMPRuntime().emitNumTeamsClause(CGF, NumTeams, ThreadLimit,
+ S.getBeginLoc());
+ }
+ };
+
+ const Expr *IfCond = nullptr;
+ for (const auto *C : S.getClausesOfKind<OMPIfClause>()) {
+ if (C->getNameModifier() == OMPD_unknown ||
+ C->getNameModifier() == OMPD_teams) {
+ IfCond = C->getCondition();
+ break;
+ }
+ }
+ if (IfCond && CGF.CGM.getLangOpts().OpenMP > 51) {
+ auto &&SerialLeague = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+ // OpenMP 6.0, 12.2, teams Construct
+ // When an if clause is present on a teams construct and the if clause
+ // expression evaluates to false, the number of formed teams is one.
+ const llvm::APInt One(32, 1);
+ auto NumTeams = IntegerLiteral(
+ CGF.getContext(), One,
+ CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
+ SourceLocation());
+ // the thread_limit clause is unaffected by the if clause
+ const auto *TL = S.getSingleClause<OMPThreadLimitClause>();
+ const Expr *ThreadLimit = TL ? TL->getThreadLimit().front() : nullptr;
+ CGF.CGM.getOpenMPRuntime().emitNumTeamsClause(CGF, &NumTeams, ThreadLimit,
+ S.getBeginLoc());
+ };
+ CGF.CGM.getOpenMPRuntime().emitIfClause(CGF, IfCond, ParallelLeague,
+ SerialLeague);
+ } else {
+ const RegionCodeGenTy ThenRCG(ParallelLeague);
+ ThenRCG(CGF);
}
OMPTeamsScope Scope(CGF, S);
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_openmp52_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_openmp52_codegen.cpp
new file mode 100644
index 0000000000000..be3118906d197
--- /dev/null
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_openmp52_codegen.cpp
@@ -0,0 +1,12098 @@
+// 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 _
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK3
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK3
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK5
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK5
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK7
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK7
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK9
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK9
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK11
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK11
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK13
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK15
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -DOMP5 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+void fn1();
+void fn2();
+void fn3();
+void fn4();
+void fn5();
+void fn6();
+void fn7();
+
+int Arg;
+
+void gtid_test() {
+#ifdef OMP5
+#pragma omp target teams distribute parallel for simd if(simd: true) nontemporal(Arg)
+#else
+#pragma omp target teams distribute parallel for simd
+#endif // OMP5
+ for (int i = 0; i < 100; i++) {
+ Arg = 0;
+ }
+
+#pragma omp target teams distribute parallel for simd if (parallel: false)
+ for(int i = 0 ; i < 100; i++) {
+ gtid_test();
+ }
+}
+
+
+void gtid_test2() {
+#ifdef OMP5
+#pragma omp target teams distribute parallel for simd if(simd: true) nontemporal(Arg)
+#else
+#pragma omp target teams distribute parallel for simd
+#endif // OMP5
+ for (int i = 0; i < 100; i++) {
+ Arg = 0;
+ }
+
+#pragma omp target teams distribute parallel for simd if (teams: false)
+ for(int i = 0 ; i < 100; i++) {
+ gtid_test2();
+ }
+}
+
+
+template <typename T>
+int tmain(T Arg) {
+#pragma omp target teams distribute parallel for simd if (true)
+ for(int i = 0 ; i < 100; i++) {
+ fn1();
+ }
+#pragma omp target teams distribute parallel for simd if (false)
+ for(int i = 0 ; i < 100; i++) {
+ fn2();
+ }
+#pragma omp target teams distribute parallel for simd if (parallel: Arg)
+ for(int i = 0 ; i < 100; i++) {
+ fn3();
+ }
+#pragma omp target teams distribute parallel for simd if (teams: Arg)
+ for(int i = 0 ; i < 100; i++) {
+ fn4();
+ }
+ return 0;
+}
+
+int main() {
+#pragma omp target teams distribute parallel for simd if (true)
+ for(int i = 0 ; i < 100; i++) {
+
+
+ fn5();
+ }
+
+#pragma omp target teams distribute parallel for simd if (false)
+ for(int i = 0 ; i < 100; i++) {
+
+
+ fn6();
+ }
+
+#pragma omp target teams distribute parallel for simd if (Arg)
+ for(int i = 0 ; i < 100; i++) {
+
+
+ fn7();
+ }
+
+ return tmain(Arg);
+}
+
+
+
+
+
+
+// call void [[T_OUTLINE_FUN_3:@.+]](
+
+#endif
+// CHECK1-LABEL: define {{[^@]+}}@_Z9gtid_testv
+// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK1-NEXT: entry:
+// CHECK1-NEXT: [[ARG_CASTED:%.*]] = alloca i64, align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
+// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4
+// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8
+// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8
+// CHECK1-NEXT: [[_TMP4:%.*]] = alloca i32, align 4
+// CHECK1-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
+// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr @Arg, align 4
+// CHECK1-NEXT: store i32 [[TMP0]], ptr [[ARG_CASTED]], align 4
+// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[ARG_CASTED]], align 8
+// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8
+// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8
+// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8
+// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8
+// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8
+// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8
+// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// CHECK1-NEXT: store i32 5, ptr [[TMP10]], align 4
+// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// CHECK1-NEXT: store i32 2, ptr [[TMP11]], align 4
+// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// CHECK1-NEXT: store ptr [[TMP8]], ptr [[TMP12]], align 8
+// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// CHECK1-NEXT: store ptr [[TMP9]], ptr [[TMP13]], align 8
+// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP14]], align 8
+// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP15]], align 8
+// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// CHECK1-NEXT: store ptr null, ptr [[TMP16]], align 8
+// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
+// CHECK1-NEXT: store ptr null, ptr [[TMP17]], align 8
+// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
+// CHECK1-NEXT: store i64 100, ptr [[TMP18]], align 8
+// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
+// CHECK1-NEXT: store i64 0, ptr [[TMP19]], align 8
+// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
+// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP20]], align 4
+// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
+// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP21]], align 4
+// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
+// CHECK1-NEXT: store i32 0, ptr [[TMP22]], align 4
+// CHECK1-NEXT: [[TMP23:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48.region_id, ptr [[KERNEL_ARGS]])
+// CHECK1-NEXT: [[TMP24:%.*]] = icmp ne i32 [[TMP23]], 0
+// CHECK1-NEXT: br i1 [[TMP24]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
+// CHECK1: omp_offload.failed:
+// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48(i64 [[TMP1]], ptr null) #[[ATTR2:[0-9]+]]
+// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]
+// CHECK1: omp_offload.cont:
+// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
+// CHECK1-NEXT: store ptr null, ptr [[TMP25]], align 8
+// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
+// CHECK1-NEXT: store ptr null, ptr [[TMP26]], align 8
+// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0
+// CHECK1-NEXT: store ptr null, ptr [[TMP27]], align 8
+// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
+// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
+// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0
+// CHECK1-NEXT: store i32 5, ptr [[TMP30]], align 4
+// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1
+// CHECK1-NEXT: store i32 1, ptr [[TMP31]], align 4
+// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2
+// CHECK1-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 8
+// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3
+// CHECK1-NEXT: store ptr [[TMP29]], ptr [[TMP33]], align 8
+// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4
+// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP34]], align 8
+// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5
+// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP35]], align 8
+// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6
+// CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8
+// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7
+// CHECK1-NEXT: store ptr null, ptr [[TMP37]], align 8
+// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8
+// CHECK1-NEXT: store i64 100, ptr [[TMP38]], align 8
+// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9
+// CHECK1-NEXT: store i64 0, ptr [[TMP39]], align 8
+// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10
+// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4
+// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11
+// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP41]], align 4
+// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12
+// CHECK1-NEXT: store i32 0, ptr [[TMP42]], align 4
+// CHECK1-NEXT: [[TMP43:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l54.region_id, ptr [[KERNEL_ARGS5]])
+// CHECK1-NEXT: [[TMP44:%.*]] = icmp ne i32 [[TMP43]], 0
+// CHECK1-NEXT: br i1 [[TMP44]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
+// CHECK1: omp_offload.failed6:
+// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l54(ptr null) #[[ATTR2]]
+// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]]
+// CHECK1: omp_offload.cont7:
+// CHECK1-NEXT: ret void
+//
+//
+// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48
+// CHECK1-SAME: (i64 noundef [[ARG:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK1-NEXT: entry:
+// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca i64, align 8
+// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
+// CHECK1-NEXT: [[ARG_CASTED:%.*]] = alloca i64, align 8
+// CHECK1-NEXT: store i64 [[ARG]], ptr [[ARG_ADDR]], align 8
+// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
+// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[ARG_ADDR]], align 4
+// CHECK1-NEXT: store i32 [[TMP0]], ptr [[ARG_CASTED]], align 4
+// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[ARG_CASTED]], align 8
+// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/207444
More information about the cfe-commits
mailing list