https://github.com/ShashwathiNavada updated
https://github.com/llvm/llvm-project/pull/141507
>From 435206a8beb2af0368757bd1cf0850db30088ca5 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshas...@pe31.hpc.amslabs.hpecorp.net>
Date: Mon, 19 May 2025 14:29:47 -0500
Subject: [PATCH 1/5] Trying to fix undefined symbol error caused by iterator
variable
---
clang/include/clang/Basic/Attr.td | 7 +++++++
clang/lib/CodeGen/CGExpr.cpp | 5 +++++
clang/lib/Sema/SemaOpenMP.cpp | 2 ++
3 files changed, 14 insertions(+)
diff --git a/clang/include/clang/Basic/Attr.td
b/clang/include/clang/Basic/Attr.td
index ccd13a4cca4dd..cc629443e9b2d 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -4461,6 +4461,13 @@ def OMPCaptureKind : Attr {
}];
}
+def OMPIterator : Attr {
+ // This attribute has no spellings as it is only ever created implicitly.
+ let Spellings = [];
+ let SemaHandler = 0;
+ let Documentation = [InternalOnly];
+}
+
def OMPReferencedVar : Attr {
// This attribute has no spellings as it is only ever created implicitly.
let Spellings = [];
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index ec01c87c13b1d..994327f690689 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2969,6 +2969,11 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction
&CGF,
llvm::Value *V = CGF.CGM.GetAddrOfGlobalVar(VD);
+ if (VD->hasAttr<OMPIteratorAttr>()) {
+ llvm::GlobalVariable *Var = llvm::dyn_cast<llvm::GlobalVariable>(V);
+ Var->setInitializer(CGF.CGM.EmitNullConstant(E->getType()));
+ }
+
if (VD->getTLSKind() != VarDecl::TLS_None)
V = CGF.Builder.CreateThreadLocalAddress(V);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index f16f841d62edd..1756c20ce486c 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -24476,6 +24476,8 @@ ExprResult SemaOpenMP::ActOnOMPIteratorExpr(Scope *S,
VarDecl::Create(Context, SemaRef.CurContext, StartLoc, D.DeclIdentLoc,
D.DeclIdent, DeclTy, TInfo, SC_None);
VD->setImplicit();
+ VD->addAttr(OMPIteratorAttr::CreateImplicit(
+ Context, SourceRange(StartLoc)));
if (S) {
// Check for conflicting previous declaration.
DeclarationNameInfo NameInfo(VD->getDeclName(), D.DeclIdentLoc);
>From 00b30ed8ff7ce3a47b81862f8e9e7e90818433e7 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshas...@pe31.hpc.amslabs.hpecorp.net>
Date: Mon, 26 May 2025 11:52:21 -0500
Subject: [PATCH 2/5] Added testcase
---
clang/test/OpenMP/declare_mapper_codegen.cpp | 102 +++++++++++++++++++
1 file changed, 102 insertions(+)
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp
b/clang/test/OpenMP/declare_mapper_codegen.cpp
index 81453223b2a27..7dc32d0ae12ff 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -1055,4 +1055,106 @@ void foo(int a){
#endif // CK4
+///==========================================================================///
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-version=52
-fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple
powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s
-o - | FileCheck --check-prefix CK5 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=52
-fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple
powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o
%t %s
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-version=52 -fopenmp
-fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown
-emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck
--check-prefix CK5 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=52
-fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple
i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=52
-fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple
powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s
-o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=52
-fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple
powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o
%t %s
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=52
-fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown
-emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck
--check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=52
-fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple
i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+
+#ifdef CK5
+typedef struct myvec {
+ int a;
+ double *b;
+} myvec_t;
+
+#pragma omp declare mapper(id: myvec_t v) map(iterator(it=0:v.a), tofrom:
v.b[it])
+// CK5: @[[ITER:[a-zA-Z0-9_]+]] = global i32 0, align 4
+
+void foo(){
+ myvec_t s;
+ #pragma omp target map(mapper(id), to:s)
+ {
+ }
+}
+
+// CK5: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*myvec[.]id]](ptr
noundef [[HANDLE:%.+]], ptr noundef [[BPTR:%.+]], ptr noundef [[BEGIN:%.+]],
i64 noundef [[BYTESIZE:%.+]], i64 noundef [[TYPE:%.+]], ptr{{.*}})
+// CK5-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], {{.*}}
+// CK5-DAG: [[PTREND:%.+]] = getelementptr %struct.myvec, ptr [[BEGIN]], i64
[[SIZE]]
+// CK5-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
+// CK5-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]
+// CK5-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
+// CK5-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
+// CK5-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
+// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
+// CK5-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK5-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK5-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
+// CK5: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK5: [[INITEVALDEL]]
+// CK5-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], {{.*}}
+
+// Remove movement mappings and mark as implicit
+// CK5-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK5-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]],
ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
+// CK5: br label %[[LHEAD:[^,]+]]
+
+// CK5: [[LHEAD]]
+// CK5: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]
+// CK5: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
+// CK5: [[LBODY]]
+// CK5: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]],
%[[LCORRECT:[^,]+]] ]
+// CK5-DAG: [[ABEGIN:%.+]] = getelementptr inbounds nuw %struct.myvec, ptr
[[PTR]], i32 0, i32 1
+// CK5-DAG: load i32, ptr @[[ITER]], align 4
+// CK5-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr
[[HANDLE]])
+// CK5-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
+// CK5-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
+// CK5-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK5-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK5-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
+// CK5-DAG: [[ALLOC]]
+// CK5-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK5-DAG: br label %[[TYEND:[^,]+]]
+// CK5-DAG: [[ALLOCELSE]]
+// CK5-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK5-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK5-DAG: [[TO]]
+// CK5-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK5-DAG: br label %[[TYEND]]
+// CK5-DAG: [[TOELSE]]
+// CK5-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK5-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK5-DAG: [[FROM]]
+// CK5-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK5-DAG: br label %[[TYEND]]
+// CK5-DAG: [[TYEND]]
+// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}})
+// CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 1
+// CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
+// CK5: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK5: [[LEXIT]]
+// CK5: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
+// CK5: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK5: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], {{.*}}
+// CK5: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
+// CK5-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], {{.*}}
+
+// Remove movement mappings and mark as implicit
+// CK5-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK5-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]],
ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
+// CK5: br label %[[DONE]]
+// CK5: [[DONE]]
+// CK5: ret void
+
+#endif // CK5
+
#endif // HEADER
>From f9029e84142fe4d4fb781f7d3d4ef1cd0fe36ede Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshas...@pe31.hpc.amslabs.hpecorp.net>
Date: Mon, 26 May 2025 12:09:22 -0500
Subject: [PATCH 3/5] Minor changes
---
clang/lib/CodeGen/CGExpr.cpp | 4 ++--
clang/lib/Sema/SemaOpenMP.cpp | 4 ++--
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 994327f690689..e6be4d639a382 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2970,8 +2970,8 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction
&CGF,
llvm::Value *V = CGF.CGM.GetAddrOfGlobalVar(VD);
if (VD->hasAttr<OMPIteratorAttr>()) {
- llvm::GlobalVariable *Var = llvm::dyn_cast<llvm::GlobalVariable>(V);
- Var->setInitializer(CGF.CGM.EmitNullConstant(E->getType()));
+ llvm::GlobalVariable *Var = llvm::dyn_cast<llvm::GlobalVariable>(V);
+ Var->setInitializer(CGF.CGM.EmitNullConstant(E->getType()));
}
if (VD->getTLSKind() != VarDecl::TLS_None)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 1756c20ce486c..7eaabae14a337 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -24476,8 +24476,8 @@ ExprResult SemaOpenMP::ActOnOMPIteratorExpr(Scope *S,
VarDecl::Create(Context, SemaRef.CurContext, StartLoc, D.DeclIdentLoc,
D.DeclIdent, DeclTy, TInfo, SC_None);
VD->setImplicit();
- VD->addAttr(OMPIteratorAttr::CreateImplicit(
- Context, SourceRange(StartLoc)));
+ VD->addAttr(
+ OMPIteratorAttr::CreateImplicit(Context, SourceRange(StartLoc)));
if (S) {
// Check for conflicting previous declaration.
DeclarationNameInfo NameInfo(VD->getDeclName(), D.DeclIdentLoc);
>From dbdbe3c4b195cac585122924a3bdcd70723ba3f1 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshas...@pe31.hpc.amslabs.hpecorp.net>
Date: Wed, 28 May 2025 08:37:44 -0500
Subject: [PATCH 4/5] Changes
---
clang/lib/CodeGen/CGExpr.cpp | 8 +++++---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 6 ++++++
2 files changed, 11 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index e6be4d639a382..d71dfe243a65b 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2968,10 +2968,12 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction
&CGF,
}
llvm::Value *V = CGF.CGM.GetAddrOfGlobalVar(VD);
-
if (VD->hasAttr<OMPIteratorAttr>()) {
- llvm::GlobalVariable *Var = llvm::dyn_cast<llvm::GlobalVariable>(V);
- Var->setInitializer(CGF.CGM.EmitNullConstant(E->getType()));
+ if (auto *GV = dyn_cast<llvm::GlobalVariable>(V)) {
+ llvm::LLVMContext &Ctx = GV->getContext();
+ llvm::MDNode *MD = llvm::MDNode::get(Ctx, llvm::MDString::get(Ctx,
"omp.iterator"));
+ GV->setMetadata("omp.iterator", MD);
+ }
}
if (VD->getTLSKind() != VarDecl::TLS_None)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 000b0c3766a46..3ee7a887400e5 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8237,6 +8237,12 @@ Expected<Function *>
OpenMPIRBuilder::emitUserDefinedMapper(
BasicBlock *FromBB = BasicBlock::Create(M.getContext(), "omp.type.from");
BasicBlock *EndBB = BasicBlock::Create(M.getContext(), "omp.type.end");
Value *IsAlloc = Builder.CreateIsNull(LeftToFrom);
+ for (GlobalVariable &GV : M.globals()) {
+ if (MDNode *MD = GV.getMetadata("omp.iterator")) {
+ auto *Zero = Constant::getNullValue(GV.getValueType());
+ GV.setInitializer(Zero);
+ }
+ }
Builder.CreateCondBr(IsAlloc, AllocBB, AllocElseBB);
// In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM.
emitBlock(AllocBB, MapperFn);
>From e65ff6bb2ab6d360f189f834367eed8d0a39185d Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshas...@pe31.hpc.amslabs.hpecorp.net>
Date: Wed, 28 May 2025 09:14:18 -0500
Subject: [PATCH 5/5] Format
---
clang/lib/CodeGen/CGExpr.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index d71dfe243a65b..ae945e7a250b4 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2971,7 +2971,8 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction
&CGF,
if (VD->hasAttr<OMPIteratorAttr>()) {
if (auto *GV = dyn_cast<llvm::GlobalVariable>(V)) {
llvm::LLVMContext &Ctx = GV->getContext();
- llvm::MDNode *MD = llvm::MDNode::get(Ctx, llvm::MDString::get(Ctx,
"omp.iterator"));
+ llvm::MDNode *MD =
+ llvm::MDNode::get(Ctx, llvm::MDString::get(Ctx, "omp.iterator"));
GV->setMetadata("omp.iterator", MD);
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
