Skip to content

Commit bd9e419

Browse files
nicebertnicebert
committed
[OpenMP] Defaultmap: fixes scalar issue, adds all variable category
Fixes issue with defaultmap where scalar isn't handled correctly for present modifier. Adds all variable cateogry introduced in OpenMP 5.2 and alters existing tests for error messages to check OpenMP 5.2 defaultmap messages.
1 parent 93d7d9b commit bd9e419

13 files changed

+220
-117
lines changed

‎clang/include/clang/Basic/OpenMPKinds.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -107,6 +107,7 @@ OPENMP_DEVICE_MODIFIER(device_num)
107107
OPENMP_DEFAULTMAP_KIND(scalar)
108108
OPENMP_DEFAULTMAP_KIND(aggregate)
109109
OPENMP_DEFAULTMAP_KIND(pointer)
110+
OPENMP_DEFAULTMAP_KIND(all)
110111

111112
// Modifiers for 'defaultmap' clause.
112113
OPENMP_DEFAULTMAP_MODIFIER(alloc)

‎clang/lib/Parse/ParseOpenMP.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3835,6 +3835,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
38353835
// Get a defaultmap modifier
38363836
unsigned Modifier = getOpenMPSimpleClauseType(
38373837
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
3838+
38383839
// Set defaultmap modifier to unknown if it is either scalar, aggregate, or
38393840
// pointer
38403841
if (Modifier < OMPC_DEFAULTMAP_MODIFIER_unknown)

‎clang/lib/Sema/SemaOpenMP.cpp

Lines changed: 14 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -822,7 +822,8 @@ class DSAStackTy {
822822
return (M == OMPC_DEFAULTMAP_MODIFIER_alloc) ||
823823
(M == OMPC_DEFAULTMAP_MODIFIER_to) ||
824824
(M == OMPC_DEFAULTMAP_MODIFIER_from) ||
825-
(M == OMPC_DEFAULTMAP_MODIFIER_tofrom);
825+
(M == OMPC_DEFAULTMAP_MODIFIER_tofrom) ||
826+
(M == OMPC_DEFAULTMAP_MODIFIER_present);
826827
}
827828
return true;
828829
}
@@ -3099,11 +3100,11 @@ ExprResult SemaOpenMP::ActOnOpenMPIdExpression(Scope *CurScope,
30993100
if (TypoCorrection Corrected =
31003101
SemaRef.CorrectTypo(Id, Sema::LookupOrdinaryName, CurScope, nullptr,
31013102
CCC, Sema::CTK_ErrorRecovery)) {
3102-
SemaRef.diagnoseTypo(Corrected,
3103-
PDiag(Lookup.empty()
3104-
? diag::err_undeclared_var_use_suggest
3105-
: diag::err_omp_expected_var_arg_suggest)
3106-
<< Id.getName());
3103+
SemaRef.diagnoseTypo(
3104+
Corrected,
3105+
SemaRef.PDiag(Lookup.empty() ? diag::err_undeclared_var_use_suggest
3106+
: diag::err_omp_expected_var_arg_suggest)
3107+
<< Id.getName());
31073108
VD = Corrected.getCorrectionDeclAs<VarDecl>();
31083109
} else {
31093110
Diag(Id.getLoc(), Lookup.empty() ? diag::err_undeclared_var_use
@@ -7591,9 +7592,9 @@ SemaOpenMP::checkOpenMPDeclareVariantFunction(SemaOpenMP::DeclGroupPtrTy DG,
75917592
PartialDiagnostic::NullDiagnostic()),
75927593
PartialDiagnosticAt(
75937594
VariantRef->getExprLoc(),
7594-
PDiag(diag::err_omp_declare_variant_doesnt_support)),
7595+
SemaRef.PDiag(diag::err_omp_declare_variant_doesnt_support)),
75957596
PartialDiagnosticAt(VariantRef->getExprLoc(),
7596-
PDiag(diag::err_omp_declare_variant_diff)
7597+
SemaRef.PDiag(diag::err_omp_declare_variant_diff)
75977598
<< FD->getLocation()),
75987599
/*TemplatesSupported=*/true, /*ConstexprSupported=*/false,
75997600
/*CLinkageMayDiffer=*/true))
@@ -21665,7 +21666,9 @@ OMPClause *SemaOpenMP::ActOnOpenMPDefaultmapClause(
2166521666
bool isDefaultmapKind = (Kind != OMPC_DEFAULTMAP_unknown) ||
2166621667
(getLangOpts().OpenMP >= 50 && KindLoc.isInvalid());
2166721668
if (!isDefaultmapKind || !isDefaultmapModifier) {
21668-
StringRef KindValue = "'scalar', 'aggregate', 'pointer'";
21669+
StringRef KindValue = getLangOpts().OpenMP < 52
21670+
? "'scalar', 'aggregate', 'pointer'"
21671+
: "'scalar', 'aggregate', 'pointer', 'all'";
2166921672
if (getLangOpts().OpenMP == 50) {
2167021673
StringRef ModifierValue = "'alloc', 'from', 'to', 'tofrom', "
2167121674
"'firstprivate', 'none', 'default'";
@@ -21709,7 +21712,7 @@ OMPClause *SemaOpenMP::ActOnOpenMPDefaultmapClause(
2170921712
return nullptr;
2171021713
}
2171121714
}
21712-
if (Kind == OMPC_DEFAULTMAP_unknown) {
21715+
if (Kind == OMPC_DEFAULTMAP_unknown || Kind == OMPC_DEFAULTMAP_all) {
2171321716
// Variable category is not specified - mark all categories.
2171421717
DSAStack->setDefaultDMAAttr(M, OMPC_DEFAULTMAP_aggregate, StartLoc);
2171521718
DSAStack->setDefaultDMAAttr(M, OMPC_DEFAULTMAP_scalar, StartLoc);
@@ -21782,7 +21785,7 @@ NamedDecl *SemaOpenMP::lookupOpenMPDeclareTargetName(
2178221785
SemaRef.CorrectTypo(Id, Sema::LookupOrdinaryName, CurScope, nullptr,
2178321786
CCC, Sema::CTK_ErrorRecovery)) {
2178421787
SemaRef.diagnoseTypo(Corrected,
21785-
PDiag(diag::err_undeclared_var_use_suggest)
21788+
SemaRef.PDiag(diag::err_undeclared_var_use_suggest)
2178621789
<< Id.getName());
2178721790
checkDeclIsAllowedInOpenMPTarget(nullptr, Corrected.getCorrectionDecl());
2178821791
return nullptr;

‎clang/test/OpenMP/target_defaultmap_messages.cpp

Lines changed: 15 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
1-
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp %s -verify=expected,omp51 -Wuninitialized -DOMP51
2-
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd %s -verify=expected,omp51 -Wuninitialized -DOMP51
1+
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp %s -fopenmp-version=52 -verify=expected,omp5x,omp52 -Wuninitialized
2+
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd %s -fopenmp-version=52 -verify=expected,omp5x,omp52 -Wuninitialized
3+
4+
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp %s -fopenmp-version=51 -verify=expected,omp5x,omp51 -Wuninitialized -DOMP51
5+
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd %s -fopenmp-version=51 -verify=expected,omp5x,omp51 -Wuninitialized -DOMP51
36

47
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=50 %s -verify=expected,omp5 -Wuninitialized -DOMP5
58
// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=50 %s -verify=expected,omp5 -Wuninitialized -DOMP5
@@ -30,23 +33,23 @@ template <class T, typename S, int N, int ST>
3033
T tmain(T argc, S **argv) {
3134
#pragma omp target defaultmap // expected-error {{expected '(' after 'defaultmap'}}
3235
foo();
33-
#pragma omp target defaultmap( // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
36+
#pragma omp target defaultmap( // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
3437
foo();
35-
#pragma omp target defaultmap() // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
38+
#pragma omp target defaultmap() // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
3639
foo();
3740
#pragma omp target defaultmap(tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
3841
foo();
39-
#pragma omp target defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
42+
#pragma omp target defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp52-error {{expected 'scalar', 'aggregate', 'pointer', 'all' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
4043
foo();
4144
#pragma omp target defaultmap(tofrom) // omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
4245
foo();
4346
#pragma omp target defaultmap(tofrom, // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
4447
foo();
45-
#pragma omp target defaultmap (scalar: // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
48+
#pragma omp target defaultmap (scalar: // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp52-error {{expected 'scalar', 'aggregate', 'pointer', 'all' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
4649
foo();
4750
#pragma omp target defaultmap(tofrom, scalar // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
4851
foo();
49-
#pragma omp target defaultmap(tofrom:scalar) defaultmap(tofrom:scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp51-error {{at most one defaultmap clause for each variable-category can appear on the directive}}
52+
#pragma omp target defaultmap(tofrom:scalar) defaultmap(tofrom:scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp5x-error {{at most one defaultmap clause for each variable-category can appear on the directive}}
5053

5154
foo();
5255

@@ -93,23 +96,23 @@ T tmain(T argc, S **argv) {
9396
int main(int argc, char **argv) {
9497
#pragma omp target defaultmap // expected-error {{expected '(' after 'defaultmap'}}
9598
foo();
96-
#pragma omp target defaultmap( // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
99+
#pragma omp target defaultmap( // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
97100
foo();
98-
#pragma omp target defaultmap() // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
101+
#pragma omp target defaultmap() // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
99102
foo();
100103
#pragma omp target defaultmap(tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
101104
foo();
102-
#pragma omp target defaultmap(tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
105+
#pragma omp target defaultmap(tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp52-error {{expected 'scalar', 'aggregate', 'pointer', 'all' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
103106
foo();
104107
#pragma omp target defaultmap(tofrom) // omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
105108
foo();
106109
#pragma omp target defaultmap(tofrom, // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
107110
foo();
108-
#pragma omp target defaultmap(scalar: // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
111+
#pragma omp target defaultmap(scalar: // omp5x-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp52-error {{expected 'scalar', 'aggregate', 'pointer', 'all' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
109112
foo();
110113
#pragma omp target defaultmap(tofrom, scalar // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
111114
foo();
112-
#pragma omp target defaultmap(tofrom: scalar) defaultmap(tofrom: scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp51-error {{at most one defaultmap clause for each variable-category can appear on the directive}}
115+
#pragma omp target defaultmap(tofrom: scalar) defaultmap(tofrom: scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp5x-error {{at most one defaultmap clause for each variable-category can appear on the directive}}
113116
foo();
114117

115118
#ifdef OMP5

‎clang/test/OpenMP/target_defaultmap_scalar_codegen.cpp

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs
2+
// RUN: %clang_cc1 -emit-llvm -o - -fopenmp \
3+
// RUN: -triple x86_64-unknown-linux-gnu %s | FileCheck %s
4+
5+
// expected-no-diagnostics
6+
#ifndef HEADER
7+
#define HEADER
8+
9+
int main() {
10+
int scalar_var = 1; // scalar
11+
int errors = 0;
12+
13+
#pragma omp target enter data map(to: scalar_var)
14+
15+
#pragma omp target map(tofrom: errors) defaultmap(present)
16+
{
17+
if(scalar_var != 1) errors++;
18+
19+
scalar_var = 7;
20+
}
21+
22+
if(scalar_var != 1) errors++;
23+
24+
#pragma omp target exit data map(delete: scalar_var)
25+
}
26+
27+
#endif
28+
// CHECK-LABEL: define {{[^@]+}}@main
29+
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
30+
// CHECK-NEXT: entry:
31+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
32+
// CHECK-NEXT: [[SCALAR_VAR:%.*]] = alloca i32, align 4
33+
// CHECK-NEXT: [[ERRORS:%.*]] = alloca i32, align 4
34+
// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4
35+
// CHECK-NEXT: store i32 1, ptr [[SCALAR_VAR]], align 4
36+
// CHECK-NEXT: store i32 0, ptr [[ERRORS]], align 4
37+
// CHECK-NEXT: call void @__omp_offloading_fd00_5567988_main_l15(ptr [[SCALAR_VAR]], ptr [[ERRORS]]) #[[ATTR2:[0-9]+]]
38+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SCALAR_VAR]], align 4
39+
// CHECK-NEXT: [[CMP:%.*]] = icmp ne i32 [[TMP0]], 1
40+
// CHECK-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
41+
// CHECK: if.then:
42+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ERRORS]], align 4
43+
// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
44+
// CHECK-NEXT: store i32 [[INC]], ptr [[ERRORS]], align 4
45+
// CHECK-NEXT: br label [[IF_END]]
46+
// CHECK: if.end:
47+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[RETVAL]], align 4
48+
// CHECK-NEXT: ret i32 [[TMP2]]
49+
//
50+
//
51+
// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_fd00_5567988_main_l15
52+
// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SCALAR_VAR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ERRORS:%.*]]) #[[ATTR1:[0-9]+]] {
53+
// CHECK-NEXT: entry:
54+
// CHECK-NEXT: [[SCALAR_VAR_ADDR:%.*]] = alloca ptr, align 8
55+
// CHECK-NEXT: [[ERRORS_ADDR:%.*]] = alloca ptr, align 8
56+
// CHECK-NEXT: store ptr [[SCALAR_VAR]], ptr [[SCALAR_VAR_ADDR]], align 8
57+
// CHECK-NEXT: store ptr [[ERRORS]], ptr [[ERRORS_ADDR]], align 8
58+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SCALAR_VAR_ADDR]], align 8
59+
// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ERRORS_ADDR]], align 8
60+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4
61+
// CHECK-NEXT: [[CMP:%.*]] = icmp ne i32 [[TMP2]], 1
62+
// CHECK-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
63+
// CHECK: if.then:
64+
// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4
65+
// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1
66+
// CHECK-NEXT: store i32 [[INC]], ptr [[TMP1]], align 4
67+
// CHECK-NEXT: br label [[IF_END]]
68+
// CHECK: if.end:
69+
// CHECK-NEXT: store i32 7, ptr [[TMP0]], align 4
70+
// CHECK-NEXT: ret void
71+
//

0 commit comments

Comments
 (0)