cchen created this revision.
cchen added a reviewer: ABataev.
Herald added subscribers: cfe-commits, guansong.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.
For this example:
int a[100];
int f2 (int i, int k)
{
for (int j = 16; j < 64; j++)
{
a[i] = j;
i += 4;
}
return i;
}
Clang will crash since it does not capture k in OpenMP outlined
function (failed assertion: "DeclRefExpr for Decl not entered in
LocalDeclMap?").
By evaluating k inside the for loop, the code can run without issue.
Therefore, my change in CGStmtOpenMP.cpp is just inserting a alloca for
k to make sure the issue is due to not capturing the variable correctly.
I think the correct way might be adding a checker in SemaOpenMP to find if
there is any step expression contain any non-constant var and add them to
the parameter of OpenMP outlined function. However, I haven't figured
out how to add the var as parameter of OpenMP outlined function
(ActOnOpenMPRegionStart
is for directive not for clause).
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D71475
Files:
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/test/OpenMP/parallel_for_linear_codegen.cpp
Index: clang/test/OpenMP/parallel_for_linear_codegen.cpp
===================================================================
--- clang/test/OpenMP/parallel_for_linear_codegen.cpp
+++ clang/test/OpenMP/parallel_for_linear_codegen.cpp
@@ -28,6 +28,19 @@
float f;
char cnt;
+int a[100];
+
+int foo (int i, int k)
+{
+#pragma omp parallel for linear (i: k + 1)
+ for (int j = 16; j < 64; j++)
+ {
+ a[i] = j;
+ i += 4;
+ }
+ return i;
+}
+
// CHECK: [[S_FLOAT_TY:%.+]] = type { float }
// CHECK: [[S_INT_TY:%.+]] = type { i32 }
// CHECK-DAG: [[F:@.+]] = global float 0.0
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -21,6 +21,7 @@
#include "clang/AST/Stmt.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/Basic/PrettyStackTrace.h"
+#include <queue>
using namespace clang;
using namespace CodeGen;
using namespace llvm::omp;
@@ -1501,6 +1502,39 @@
if (const auto *CS = cast_or_null<BinaryOperator>(C->getCalcStep()))
if (const auto *SaveRef = cast<DeclRefExpr>(CS->getLHS())) {
EmitVarDecl(*cast<VarDecl>(SaveRef->getDecl()));
+ // Run DFS to find all the non-constant node under linear-step
expression
+ const Expr* Step = CS->getRHS();
+ std::queue<const Stmt*> StmtQueue;
+ StmtQueue.push(Step);
+ while (!StmtQueue.empty()) {
+ const Stmt* CurStep = StmtQueue.front();
+ StmtQueue.pop();
+ if (const auto *BO = dyn_cast<BinaryOperator>(CurStep)) {
+ StmtQueue.push(BO->getLHS());
+ StmtQueue.push(BO->getRHS());
+ } else if (isa<ImplicitCastExpr>(CurStep)) {
+ const auto *IC = dyn_cast<ImplicitCastExpr>(CurStep);
+ for (const Stmt *Child: IC->children()) {
+ StmtQueue.push(Child);
+ }
+ } else if (isa<DeclRefExpr>(CurStep)) {
+ // Generate a `alloca` for the variable in linear-step that
+ // is not inside LocalDeclMap since compiler does not capture
+ // this variable inside openmp-construct AST (which also
+ // don't consider this variable as LValue either).
+ // TODO The runtime behavior is incorrect for now since I only add
+ // an `alloca`, and haven't store a meaningful value yet.
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(CurStep)) {
+ const ValueDecl* VD = DRE->getDecl();
+ const VarDecl* VarD = cast<VarDecl>(VD);
+ if (LocalDeclMap.find(VarD) == LocalDeclMap.end()) {
+ AutoVarEmission Emission = EmitAutoVarAlloca(*VarD);
+ LocalDeclMap.insert({VarD, Emission.getAllocatedAddress()});
+ EmitAutoVarCleanups(Emission);
+ }
+ }
+ }
+ }
// Emit calculation of the linear step.
EmitIgnoredExpr(CS);
}
Index: clang/test/OpenMP/parallel_for_linear_codegen.cpp
===================================================================
--- clang/test/OpenMP/parallel_for_linear_codegen.cpp
+++ clang/test/OpenMP/parallel_for_linear_codegen.cpp
@@ -28,6 +28,19 @@
float f;
char cnt;
+int a[100];
+
+int foo (int i, int k)
+{
+#pragma omp parallel for linear (i: k + 1)
+ for (int j = 16; j < 64; j++)
+ {
+ a[i] = j;
+ i += 4;
+ }
+ return i;
+}
+
// CHECK: [[S_FLOAT_TY:%.+]] = type { float }
// CHECK: [[S_INT_TY:%.+]] = type { i32 }
// CHECK-DAG: [[F:@.+]] = global float 0.0
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -21,6 +21,7 @@
#include "clang/AST/Stmt.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/Basic/PrettyStackTrace.h"
+#include <queue>
using namespace clang;
using namespace CodeGen;
using namespace llvm::omp;
@@ -1501,6 +1502,39 @@
if (const auto *CS = cast_or_null<BinaryOperator>(C->getCalcStep()))
if (const auto *SaveRef = cast<DeclRefExpr>(CS->getLHS())) {
EmitVarDecl(*cast<VarDecl>(SaveRef->getDecl()));
+ // Run DFS to find all the non-constant node under linear-step expression
+ const Expr* Step = CS->getRHS();
+ std::queue<const Stmt*> StmtQueue;
+ StmtQueue.push(Step);
+ while (!StmtQueue.empty()) {
+ const Stmt* CurStep = StmtQueue.front();
+ StmtQueue.pop();
+ if (const auto *BO = dyn_cast<BinaryOperator>(CurStep)) {
+ StmtQueue.push(BO->getLHS());
+ StmtQueue.push(BO->getRHS());
+ } else if (isa<ImplicitCastExpr>(CurStep)) {
+ const auto *IC = dyn_cast<ImplicitCastExpr>(CurStep);
+ for (const Stmt *Child: IC->children()) {
+ StmtQueue.push(Child);
+ }
+ } else if (isa<DeclRefExpr>(CurStep)) {
+ // Generate a `alloca` for the variable in linear-step that
+ // is not inside LocalDeclMap since compiler does not capture
+ // this variable inside openmp-construct AST (which also
+ // don't consider this variable as LValue either).
+ // TODO The runtime behavior is incorrect for now since I only add
+ // an `alloca`, and haven't store a meaningful value yet.
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(CurStep)) {
+ const ValueDecl* VD = DRE->getDecl();
+ const VarDecl* VarD = cast<VarDecl>(VD);
+ if (LocalDeclMap.find(VarD) == LocalDeclMap.end()) {
+ AutoVarEmission Emission = EmitAutoVarAlloca(*VarD);
+ LocalDeclMap.insert({VarD, Emission.getAllocatedAddress()});
+ EmitAutoVarCleanups(Emission);
+ }
+ }
+ }
+ }
// Emit calculation of the linear step.
EmitIgnoredExpr(CS);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
