jdenny created this revision.
jdenny added reviewers: jhuber6, jdoerfert.
Herald added subscribers: guansong, yaxunl.
jdenny requested review of this revision.
Herald added a subscriber: sstefan1.
Herald added a project: clang.
This patch fixes a bug from D89802 <https://reviews.llvm.org/D89802>. For
example, without it, Clang
generates x as the debug map name for both x and y in the following
example:
#pragma omp target map(to: x, y)
x = y = 1;
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D101564
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_map_names.cpp
Index: clang/test/OpenMP/target_map_names.cpp
===================================================================
--- clang/test/OpenMP/target_map_names.cpp
+++ clang/test/OpenMP/target_map_names.cpp
@@ -16,7 +16,6 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
-// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
@@ -25,7 +24,6 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
-// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";ps->p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
@@ -181,6 +179,18 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// Clang used to mistakenly generate the map name "x" for both x and y on this
+// directive. Conditions to reproduce the bug: a single map clause has two
+// variables, and at least the second is used in the associated statement.
+//
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";x;{{.*}}.cpp;[[@LINE+3]];7;;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8]
c";y;{{.*}}.cpp;[[@LINE+2]];10;;\00"
+void secondMapNameInClause() {
+ int x, y;
+ #pragma omp target map(to: x, y)
+ x = y = 1;
+}
+
// DEBUG: %{{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64*
{{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]*
@.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
// DEBUG: %{{.+}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t*
@{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}},
i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x
i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}, i32 {{.+}}, i32
{{.+}})
// DEBUG: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}},
i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8**
getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]*
@.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7764,6 +7764,8 @@
const ValueDecl *MapDecl = (I->getAssociatedDeclaration())
? I->getAssociatedDeclaration()
: BaseDecl;
+ MapExpr = (I->getAssociatedExpression()) ? I->getAssociatedExpression()
+ : MapExpr;
// Get information on whether the element is a pointer. Have to do a
// special treatment for array sections given that they are built-in
Index: clang/test/OpenMP/target_map_names.cpp
===================================================================
--- clang/test/OpenMP/target_map_names.cpp
+++ clang/test/OpenMP/target_map_names.cpp
@@ -16,7 +16,6 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
-// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
@@ -25,7 +24,6 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
-// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
@@ -181,6 +179,18 @@
// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// Clang used to mistakenly generate the map name "x" for both x and y on this
+// directive. Conditions to reproduce the bug: a single map clause has two
+// variables, and at least the second is used in the associated statement.
+//
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;[[@LINE+3]];7;;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";y;{{.*}}.cpp;[[@LINE+2]];10;;\00"
+void secondMapNameInClause() {
+ int x, y;
+ #pragma omp target map(to: x, y)
+ x = y = 1;
+}
+
// DEBUG: %{{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
// DEBUG: %{{.+}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}, i32 {{.+}}, i32 {{.+}})
// DEBUG: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7764,6 +7764,8 @@
const ValueDecl *MapDecl = (I->getAssociatedDeclaration())
? I->getAssociatedDeclaration()
: BaseDecl;
+ MapExpr = (I->getAssociatedExpression()) ? I->getAssociatedExpression()
+ : MapExpr;
// Get information on whether the element is a pointer. Have to do a
// special treatment for array sections given that they are built-in
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
