Merge branch 'sotoc-issue-44' into 'aurora_offloading_prototype'

Resolve sotoc issue 44

Closes llvm#44

See merge request NEC-RWTH-Projects/clang!24

Author: manorom

clang/tools/sotoc/src/DeclResolver.cpp

@@ -182,7 +182,7 @@ void TypeDeclResolver::runOwnVisitor(clang::Decl *D,
 
 void FunctionDeclResolver::runOwnVisitor(
     clang::Decl *D, std::function<void(clang::Decl *Dep)> Fn) {
-  DEBUGPDECL(D, "Searching for referred decls in function " );
+  DEBUGPDECL(D, "Searching for referred decls in function ");
   DiscoverFunctionsInDeclVisitor Visitor(Fn);
   Visitor.TraverseDecl(D);
 }

clang/tools/sotoc/src/TargetCode.cpp

@@ -125,10 +125,12 @@ void TargetCode::generateFunctionPrologue(TargetCodeRegion *TCR,
       int dim = 0;
 
       std::vector<int> VariableDimensions;
-      handleArrays(&t, DimString, dim, VariableDimensions, TCR, elemType, VarName);
+      handleArrays(&t, DimString, dim, VariableDimensions, TCR, elemType,
+                   VarName);
 
       for (int d : VariableDimensions) {
-        Out << "unsigned long long __sotoc_vla_dim" << d << "_" << VarName << ", ";
+        Out << "unsigned long long __sotoc_vla_dim" << d << "_" << VarName
+            << ", ";
       }
 
       // set type to void* to avoid warnings from the compiler
@@ -178,6 +180,13 @@ void TargetCode::generateFunctionPrologue(TargetCodeRegion *TCR,
 
       Out << " = __sotoc_var_" << VarName << ";\n";
 
+      auto LowerBound = TCR->CapturedLowerBounds.find(*I);
+      if (LowerBound != TCR->CapturedLowerBounds.end()) {
+        Out << VarName << " = " << VarName << " - (";
+        LowerBound->second->printPretty(Out, NULL, TCR->getPP());
+        Out << ");\n";
+      }
+
     } else {
       if (!(*I)->getType().getTypePtr()->isPointerType()) {
         if (C) {
@@ -267,8 +276,7 @@ std::string TargetCode::generateFunctionName(TargetCodeRegion *TCR) {
 void TargetCode::handleArrays(const clang::ArrayType **t,
                               std::list<std::string> &DimString, int &dim,
                               std::vector<int> &VariableDims,
-                              TargetCodeRegion *TCR,
-                              std::string &elemType,
+                              TargetCodeRegion *TCR, std::string &elemType,
                               const std::string &ArrayName) {
   auto OrigT = *t;
 

clang/tools/sotoc/src/TargetCode.h

@@ -85,7 +85,6 @@ class TargetCode {
   /// \param returns the last element type (i.e., the type of the array)
   void handleArrays(const clang::ArrayType **t,
                     std::list<std::string> &DimString, int &dim,
-                    std::vector<int> &VariableDimensions,
-                    TargetCodeRegion *TCR, std::string &elemType,
-                    const std::string &ArrayName);
+                    std::vector<int> &VariableDimensions, TargetCodeRegion *TCR,
+                    std::string &elemType, const std::string &ArrayName);
 };

clang/tools/sotoc/src/TargetCodeFragment.h

@@ -159,6 +159,12 @@ class TargetCodeRegion : public TargetCodeFragment {
   /// location of the first pragma of the target region to compose the name of
   ///  the function generated for that region)
   clang::SourceLocation getTargetDirectiveLocation();
+  /// Lower bounds of mapped array slices (if lower then 0).
+  /// If the captured variable is an array, of which only a slice is mapped
+  /// (by a map() clause), the incoming pointer argument will need to be
+  /// shifted to the right if the lower bound of that slice is not 0.
+  /// If this is the case, the lower bound is saved into this map.
+  std::map<clang::VarDecl *, clang::Expr *> CapturedLowerBounds;
 };
 
 /// This class represents a declaration, i.e. a function, global varialbe, or

clang/tools/sotoc/src/Visitors.cpp

@@ -19,6 +19,7 @@
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/Decl.h"
+#include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Basic/OpenMPKinds.h"
@@ -119,11 +120,12 @@ bool FindTargetCodeVisitor::VisitStmt(clang::Stmt *S) {
   return true;
 }
 
-class CollectOMPClausesVisitor : public clang::RecursiveASTVisitor<CollectOMPClausesVisitor> {
+class CollectOMPClausesVisitor
+    : public clang::RecursiveASTVisitor<CollectOMPClausesVisitor> {
   std::shared_ptr<TargetCodeRegion> TCR;
+
 public:
-  CollectOMPClausesVisitor(std::shared_ptr<TargetCodeRegion> &TCR)
-      :  TCR(TCR) {};
+  CollectOMPClausesVisitor(std::shared_ptr<TargetCodeRegion> &TCR) : TCR(TCR){};
   bool VisitStmt(clang::Stmt *S) {
     if (auto *OED = llvm::dyn_cast<clang::OMPExecutableDirective>(S)) {
       for (auto *Clause : OED->clauses()) {
@@ -153,6 +155,9 @@ bool FindTargetCodeVisitor::processTargetRegion(
       // if the target region cannot be added we dont want to parse its args
       if (TargetCodeInfo.addCodeFragment(TCR)) {
 
+        FindArraySectionVisitor ArraySectionVisitor(TCR->CapturedLowerBounds);
+        ArraySectionVisitor.TraverseStmt(TargetDirective);
+
         // look for nested clause
         CollectOMPClausesVisitor(TCR).TraverseStmt(CS);
         for (auto C : TargetDirective->clauses()) {
@@ -198,7 +203,7 @@ void FindTargetCodeVisitor::addTargetRegionArgs(
     DEBUGP("Iterating var set");
     // i->print(llvm::outs());
     if (Context.getSourceManager().isBeforeInTranslationUnit(
-          S->getBeginLoc(),i->getSourceRange().getBegin())) {
+            S->getBeginLoc(), i->getSourceRange().getBegin())) {
       tmpSet.insert(i);
       continue;
     }
@@ -365,3 +370,35 @@ bool DiscoverFunctionsInDeclVisitor::VisitExpr(clang::Expr *E) {
   }
   return true;
 }
+
+bool FindArraySectionVisitor::VisitExpr(clang::Expr *E) {
+  if (auto *ASE = llvm::dyn_cast<clang::OMPArraySectionExpr>(E)) {
+    clang::Expr *Base = ASE->getBase();
+    if (llvm::isa<clang::OMPArraySectionExpr>(Base)) {
+      return true;
+    }
+    if (auto *CastBase = llvm::dyn_cast<clang::CastExpr>(Base)) {
+      Base = CastBase->getSubExpr();
+      if (auto *DRE = llvm::dyn_cast<clang::DeclRefExpr>(Base)) {
+        auto *VarDecl = llvm::dyn_cast<clang::VarDecl>(DRE->getDecl());
+        if (!VarDecl) {
+          llvm::errs() << "VALDECL != VARDECL\n";
+          return true;
+        }
+        clang::Expr *LowerBound = ASE->getLowerBound();
+        if (!LowerBound) {
+          return true;
+        }
+
+        if (auto *IntegerLiteral =
+                llvm::dyn_cast<clang::IntegerLiteral>(LowerBound)) {
+          if (IntegerLiteral->getValue() == 0) {
+            return true;
+          }
+        }
+        LowerBoundsMap.emplace(VarDecl, LowerBound);
+      }
+    }
+  }
+  return true;
+}

clang/tools/sotoc/src/Visitors.h

@@ -149,3 +149,15 @@ class FindTargetCodeVisitor
   void addTargetRegionArgs(clang::CapturedStmt *S,
                            std::shared_ptr<TargetCodeRegion> TCR);
 };
+
+class FindArraySectionVisitor
+    : public clang::RecursiveASTVisitor<FindArraySectionVisitor> {
+
+  std::map<clang::VarDecl *, clang::Expr *> &LowerBoundsMap;
+
+public:
+  FindArraySectionVisitor(
+      std::map<clang::VarDecl *, clang::Expr *> &LowerBoundsMap)
+      : LowerBoundsMap(LowerBoundsMap) {}
+  bool VisitExpr(clang::Expr *E);
+};

clang/tools/sotoc/test/arrays/slice_mixed_length_array_2d.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){

clang/tools/sotoc/test/arrays/slice_static_length_array_1d.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){

clang/tools/sotoc/test/arrays/slice_variable_length_array_1d.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){

clang/tools/sotoc/test/arrays/slice_variable_length_array_2d.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){
@@ -8,7 +8,7 @@ int main(){
   int sizeY=512;
   float A[sizeX][sizeY]; 
 
-  #pragma omp target map(tofrom:A[:sizeX/2][:sizeY])
+  #pragma omp target map(tofrom:A[0:sizeX/2][:sizeY])
   {
     int i;
     int j;

clang/tools/sotoc/test/arrays/slice_variable_length_array_2d_1_colon.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){

clang/tools/sotoc/test/arrays/slice_variable_length_array_2d_1_explicit.c

@@ -1,5 +1,5 @@
 // RUN: %sotoc-transform-compile
-// #RUN#: %run-on-host | %filecheck %s
+// RUN: %run-on-host | %filecheck %s
 
 #include <stdio.h>
 int main(){

clang/tools/sotoc/test/arrays/slice_variable_length_array_3d.c

@@ -15,7 +15,7 @@ int main(){
     int j;
     int k;
     for(i=0; i< 512; i++){
-      for(j=0 ; j< 256; j+=2){
+      for(j=0 ; j< 512; j+=1){
         for(k=0; k < (sizeZ - 2); ++k) {
           A[i][j][k]=i-j;
         }