https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/162499
When an array section bound was in a list of bound-types we improperly
generated the list of types of the bounds because
getPointeeOrElementType gets the LOWEST level of type (that is, digs
through ALL array types to get to the base-est of types) when what
we really wanted was 1 layer of pointer/array removed.
This patch fixes it and adds a test that showed the problem by re-ordering the
existing ones. This wasn't previously obvious by chance, since the
'array-index-only' variants ended up generating the recipe, and not the bounds.
>From a3e4c2ff2c7d06130fb4b87bec4257160105a8ca Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Wed, 8 Oct 2025 07:52:51 -0700
Subject: [PATCH] [OpenACC] Fix issue with array-section type generation
When an array section bound was in a list of bound-types we improperly
generated the list of types of the bounds because
getPointeeOrElementType gets the LOWEST level of type (that is, digs
through ALL array types to get to the base-est of types) when what
we really wanted was 1 layer of pointer/array removed.
This patch fixes it and adds a test that showed the problem by
re-ordering the existing ones. This wasn't previously obvious by
chance, since the 'array-index-only' variants ended up generating the
recipe, and not the bounds.
---
clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp | 5 ++++-
.../CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp | 4 ++--
2 files changed, 6 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index a9af753381db3..4cf2237468afd 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -87,7 +87,10 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
if (const auto *section = dyn_cast<ArraySectionExpr>(curVarExpr)) {
QualType baseTy = ArraySectionExpr::getBaseOriginalType(
section->getBase()->IgnoreParenImpCasts());
- boundTypes.push_back(QualType(baseTy->getPointeeOrArrayElementType(),
0));
+ if (auto *at = getContext().getAsArrayType(baseTy))
+ boundTypes.push_back(at->getElementType());
+ else
+ boundTypes.push_back(baseTy->getPointeeType());
} else {
boundTypes.push_back(curVarExpr->getType());
}
diff --git a/clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp
b/clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp
index e83e548cd138b..74cb567076f21 100644
--- a/clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp
@@ -21,7 +21,7 @@ void do_things(unsigned A, unsigned B) {
;
T TwoArr[5][5];
-#pragma acc parallel private(TwoArr[B][B])
+#pragma acc parallel private(TwoArr[A:B][A:B])
// CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSA5_A5_i :
!cir.ptr<!cir.array<!cir.array<!s32i x 5> x 5>> init {
// CHECK-NEXT: ^bb0(%arg0: !cir.ptr<!cir.array<!cir.array<!s32i x 5> x 5>>
{{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]:
!acc.data_bounds_ty {{.*}}):
// CHECK-NEXT: %[[TL_ALLOCA:.*]] = cir.alloca !cir.array<!cir.array<!s32i x 5>
x 5>, !cir.ptr<!cir.array<!cir.array<!s32i x 5> x 5>>, ["openacc.private.init"]
{alignment = 4 : i64}
@@ -30,7 +30,7 @@ void do_things(unsigned A, unsigned B) {
;
#pragma acc parallel private(TwoArr[B][A:B])
;
-#pragma acc parallel private(TwoArr[A:B][A:B])
+#pragma acc parallel private(TwoArr[B][B])
;
#pragma acc parallel private(TwoArr)
// CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_A5_i :
!cir.ptr<!cir.array<!cir.array<!s32i x 5> x 5>> init {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
