https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/138196
Previously we just checked duplicates for a handful of clauses between active
device_type clauses. However, after looking into the implementation for
'collapse', I discovered that duplicate device_type identifiers with duplicate
versions of htese clause are problematic.
This patch corrects this and now catches these conflicts.
>From 3b0bd9a11e85324d2b00566a6816b03930afcfe7 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Thu, 1 May 2025 13:14:51 -0700
Subject: [PATCH] [OpenACC] Implement better dupe catching for device_type
clauses
Previously we just checked duplicates for a handful of clauses between
active device_type clauses. However, after looking into the
implementation for 'collapse', I discovered that duplicate device_type
identifiers with duplicate versions of htese clause are problematic.
This patch corrects this and now catches these conflicts.
---
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/lib/Sema/SemaOpenACCClause.cpp | 98 ++++++++++++++++---
.../compute-construct-num_gangs-clause.c | 15 +++
.../compute-construct-num_workers-clause.c | 15 +++
.../compute-construct-vector_length-clause.c | 15 +++
.../loop-construct-collapse-clause.cpp | 21 ++++
.../loop-construct-tile-clause.cpp | 16 +++
7 files changed, 171 insertions(+), 12 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 604bb56d28252..f279d84136562 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13075,6 +13075,9 @@ def err_acc_clause_routine_one_of_in_region
def err_acc_clause_since_last_device_type
: Error<"OpenACC '%0' clause cannot appear more than once%select{| in a "
"'device_type' region}2 on a '%1' directive">;
+def err_acc_clause_conflicts_prev_dev_type
+ : Error<"OpenACC '%0' clause applies to 'device_type' '%1', which "
+ "conflicts with previous clause">;
def err_acc_reduction_num_gangs_conflict
: Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not "
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp
b/clang/lib/Sema/SemaOpenACCClause.cpp
index 6cf6888e2a3a9..87dae457ac3e9 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -335,29 +335,103 @@ class SemaOpenACCClauseVisitor {
// For 'tile' and 'collapse', only allow 1 per 'device_type'.
// Also applies to num_worker, num_gangs, vector_length, and async.
+ // This does introspection into the actual device-types to prevent duplicates
+ // across device types as well.
template <typename TheClauseTy>
bool DisallowSinceLastDeviceType(SemaOpenACC::OpenACCParsedClause &Clause) {
auto LastDeviceTypeItr =
std::find_if(ExistingClauses.rbegin(), ExistingClauses.rend(),
llvm::IsaPred<OpenACCDeviceTypeClause>);
- auto Last = std::find_if(ExistingClauses.rbegin(), LastDeviceTypeItr,
- llvm::IsaPred<TheClauseTy>);
+ auto LastSinceDevTy =
+ std::find_if(ExistingClauses.rbegin(), LastDeviceTypeItr,
+ llvm::IsaPred<TheClauseTy>);
- if (Last == LastDeviceTypeItr)
+ // In this case there is a duplicate since the last device_type/lack of a
+ // device_type. Diagnose these as duplicates.
+ if (LastSinceDevTy != LastDeviceTypeItr) {
+ SemaRef.Diag(Clause.getBeginLoc(),
+ diag::err_acc_clause_since_last_device_type)
+ << Clause.getClauseKind() << Clause.getDirectiveKind()
+ << (LastDeviceTypeItr != ExistingClauses.rend());
+ SemaRef.Diag((*LastSinceDevTy)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+
+ // Mention the last device_type as well.
+ if (LastDeviceTypeItr != ExistingClauses.rend())
+ SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ return true;
+ }
+
+ // If this isn't in a device_type, and we didn't diagnose that there are
+ // dupes above, just give up, no sense in searching for previous
device_type
+ // regions as they don't exist.
+ if (LastDeviceTypeItr == ExistingClauses.rend())
return false;
- SemaRef.Diag(Clause.getBeginLoc(),
- diag::err_acc_clause_since_last_device_type)
- << Clause.getClauseKind() << Clause.getDirectiveKind()
- << (LastDeviceTypeItr != ExistingClauses.rend());
- SemaRef.Diag((*Last)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ // The device-type that is active for us, so we can compare to the previous
+ // ones.
+ const auto &ActiveDeviceTypeClause =
+ cast<OpenACCDeviceTypeClause>(**LastDeviceTypeItr);
- if (LastDeviceTypeItr != ExistingClauses.rend())
- SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
- diag::note_acc_previous_clause_here);
+ auto PrevDeviceTypeItr = LastDeviceTypeItr;
+ auto CurDevTypeItr = LastDeviceTypeItr;
- return true;
+ while ((CurDevTypeItr = std::find_if(
+ std::next(PrevDeviceTypeItr), ExistingClauses.rend(),
+ llvm::IsaPred<OpenACCDeviceTypeClause>)) !=
+ ExistingClauses.rend()) {
+ // At this point, we know that we have a region between two device_types,
+ // as specified by CurDevTypeItr and PrevDeviceTypeItr.
+
+ auto CurClauseKindItr = std::find_if(PrevDeviceTypeItr, CurDevTypeItr,
+ llvm::IsaPred<TheClauseTy>);
+
+ // There are no clauses of the current kind between these device_types,
so
+ // continue.
+ if (CurClauseKindItr == CurDevTypeItr)
+ continue;
+
+ // At this point, we know that this device_type region has a collapse.
So
+ // diagnose if the two device_types have any overlap in their
+ // architectures.
+ const auto &CurDeviceTypeClause =
+ cast<OpenACCDeviceTypeClause>(**CurDevTypeItr);
+
+ for (const DeviceTypeArgument &arg :
+ ActiveDeviceTypeClause.getArchitectures()) {
+ for (const DeviceTypeArgument &prevArg :
+ CurDeviceTypeClause.getArchitectures()) {
+
+ // This should catch duplicates * regions, duplicate same-text
(thanks
+ // to identifier equiv.) and case insensitive dupes.
+ if (arg.getIdentifierInfo() == prevArg.getIdentifierInfo() ||
+ (arg.getIdentifierInfo() && prevArg.getIdentifierInfo() &&
+
StringRef{arg.getIdentifierInfo()->getName()}.equals_insensitive(
+ prevArg.getIdentifierInfo()->getName()))) {
+ SemaRef.Diag(Clause.getBeginLoc(),
+ diag::err_acc_clause_conflicts_prev_dev_type)
+ << Clause.getClauseKind()
+ << (arg.getIdentifierInfo() ?
arg.getIdentifierInfo()->getName()
+ : "*");
+ // mention the active device type.
+ SemaRef.Diag(ActiveDeviceTypeClause.getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ // mention the previous clause.
+ SemaRef.Diag((*CurClauseKindItr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ // mention the previous device type.
+ SemaRef.Diag(CurDeviceTypeClause.getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ return true;
+ }
+ }
+ }
+
+ PrevDeviceTypeItr = CurDevTypeItr;
+ }
+ return false;
}
public:
diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
index 9651514ca0fff..9db6f461081c6 100644
--- a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
@@ -71,3 +71,18 @@ void Test() {
#pragma acc loop num_gangs(1)
for(int i = 5; i < 10;++i);
}
+
+void no_dupes_since_last_device_type() {
+ // expected-error@+4{{OpenACC 'num_gangs' clause applies to 'device_type'
'radeon', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) num_gangs(getS())
device_type(radeon) num_gangs(getS())
+ ;
+ // expected-error@+4{{OpenACC 'num_gangs' clause applies to 'device_type'
'nvidia', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) num_gangs(getS()) device_type(nvidia,
radeon) num_gangs(getS())
+ ;
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
index 23a73f0bf49bd..937d15cb1b65b 100644
--- a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
@@ -60,3 +60,18 @@ void Test() {
#pragma acc loop num_workers(1)
for(int i = 5; i < 10;++i);
}
+
+void no_dupes_since_last_device_type() {
+ // expected-error@+4{{OpenACC 'num_workers' clause applies to 'device_type'
'radEon', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) num_workers(getS())
device_type(radEon) num_workers(getS())
+ ;
+ // expected-error@+4{{OpenACC 'num_workers' clause applies to 'device_type'
'nvidia', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) num_workers(getS())
device_type(nvidia, radeon) num_workers(getS())
+ ;
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
index c79fc627252f1..e0d7571af687f 100644
--- a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
@@ -60,3 +60,18 @@ void Test() {
#pragma acc loop vector_length(1)
for(int i = 5; i < 10;++i);
}
+
+void no_dupes_since_last_device_type() {
+ // expected-error@+4{{OpenACC 'vector_length' clause applies to
'device_type' 'radeon', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) vector_length(getS())
device_type(radeon) vector_length(getS())
+ ;
+ // expected-error@+4{{OpenACC 'vector_length' clause applies to
'device_type' 'nvidia', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) vector_length(getS())
device_type(nvidia, radeon) vector_length(getS())
+ ;
+}
diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
index 851742434a1f9..ea100d8290a19 100644
--- a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
@@ -539,4 +539,25 @@ void no_dupes_since_last_device_type() {
#pragma acc loop collapse(1) device_type(*) collapse(1) device_type(nvidia)
collapse(2)
for(unsigned i = 0; i < 5; ++i)
for(unsigned j = 0; j < 5; ++j);
+
+ // This one is ok, despite * being the 'all' value.
+#pragma acc loop device_type(*) collapse(1) device_type(nvidia) collapse(2)
+ for(unsigned i = 0; i < 5; ++i)
+ for(unsigned j = 0; j < 5; ++j);
+
+ // expected-error@+4{{OpenACC 'collapse' clause applies to 'device_type'
'nvidia', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(nviDia, radeon) collapse(1) device_type(nvidia)
collapse(2)
+ for(unsigned i = 0; i < 5; ++i)
+ for(unsigned j = 0; j < 5; ++j);
+
+ // expected-error@+4{{OpenACC 'collapse' clause applies to 'device_type'
'radeon', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(radeon) collapse(1) device_type(nvidia, radeon)
collapse(2)
+ for(unsigned i = 0; i < 5; ++i)
+ for(unsigned j = 0; j < 5; ++j);
}
diff --git a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
index ebd657791f34d..3fe7e0582b77e 100644
--- a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
@@ -419,4 +419,20 @@ void no_dupes_since_last_device_type() {
#pragma acc loop tile(1) device_type(*) tile(1) device_type(nvidia) tile(2)
for(unsigned i = 0; i < 5; ++i)
for(unsigned j = 0; j < 5; ++j);
+
+ // expected-error@+4{{OpenACC 'tile' clause applies to 'device_type'
'nvidiA', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(nvidia, radeon) tile(1) device_type(nvidiA)
tile(2)
+ for(unsigned i = 0; i < 5; ++i)
+ for(unsigned j = 0; j < 5; ++j);
+
+ // expected-error@+4{{OpenACC 'tile' clause applies to 'device_type'
'radeon', which conflicts with previous clause}}
+ // expected-note@+3{{previous clause is here}}
+ // expected-note@+2{{previous clause is here}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(radeon) tile(1) device_type(nvidia, radeon)
tile(2)
+ for(unsigned i = 0; i < 5; ++i)
+ for(unsigned j = 0; j < 5; ++j);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
