Merge branch 'sotoc-vla-arguments' into 'aurora_offloading_prototype'

Fix generated target region parameters for multidimensional VLAs Closes llvm#28 See merge request NEC-RWTH-Projects/clang!21
RWTH-HPC · Apr 18, 2019 · bcb9973 · bcb9973
2 parents 2acd695 + 219260b
commit bcb9973
Show file tree

Hide file tree

Showing 13 changed files with 115 additions and 28 deletions.
diff --git a/clang/tools/sotoc/CMakeLists.txt b/clang/tools/sotoc/CMakeLists.txt
@@ -57,22 +57,8 @@ if (SOTOC_STANDALONE_BUILD)
 ################################################################################
 
 else()
-  #set(CLANG_CONFIG_HINT ${CMAKE_BINARY_DIR}/lib/cmake/clang)
-  #set(CLANG_LIBRARY_DIRS ${CMAKE_BINARY_DIR}/lib)
-  #set(CLANG_INCLUDE_DIRS ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/tools/clang/include )
-  #libomptarget_say("_________A__________${CLANG_INCLUDE_DIRS}")
-  #libomptarget_say("_________B__________${CLANG_LIBRARY_DIRS}")
-  #libomptarget_say("_________C__________${LLVMCONFIG_FILE}")
   add_clang_executable(${SOTOC_PROJECT_NAME} ${SOTOC_SOURCES})
 endif()
-#  include_directories(${CLANG_INCLUDE_DIRS})
-#  link_directories(${CLANG_LIBRARY_DIRS})
-#  add_definitions(${CLANG_DEFINITIONS})
-
-#  add_definitions(
-#  -D__STDC_LIMIT_MACROS
-#  -D__STDC_CONSTANT_MACROS
-#  )
 
 set(SOTOC_DEBUG_OUTPUT OFF CACHE BOOL "Compiles sotoc with possible debug output")
 if (SOTOC_DEBUG_OUTPUT)

diff --git a/clang/tools/sotoc/src/TargetCode.cpp b/clang/tools/sotoc/src/TargetCode.cpp
@@ -77,7 +77,7 @@ void TargetCode::generateCode(llvm::raw_ostream &Out) {
     if (TCR) {
       generateFunctionPrologue(TCR, Out);
     }
-    
+
     Out << Frag->PrintPretty();
 
     if (TCR) {
@@ -103,6 +103,7 @@ void TargetCode::generateFunctionPrologue(TargetCodeRegion *TCR,
   Out << "void " << generateFunctionName(TCR) << "(";
   for (auto i = TCR->getCapturedVarsBegin(), e = TCR->getCapturedVarsEnd();
        i != e; ++i) {
+    std::string VarName = (*i)->getDeclName().getAsString();
     auto C = TCR->GetReferredOMPClause(*i);
     if (!first) {
       Out << ", ";
@@ -116,7 +117,12 @@ void TargetCode::generateFunctionPrologue(TargetCodeRegion *TCR,
       DEBUGP("Generating code for array type");
       int dim = 0;
 
-      handleArrays(&t, DimString, dim, TCR, elemType);
+      std::vector<int> VariableDimensions;
+      handleArrays(&t, DimString, dim, VariableDimensions, TCR, elemType, VarName);
+
+      for (int d : VariableDimensions) {
+        Out << "unsigned long long __sotoc_vla_dim" << d << "_" << VarName << ", ";
+      }
 
       // set type to void* to avoid warnings from the compiler
       Out << "void *__sotoc_var_";
@@ -135,7 +141,7 @@ void TargetCode::generateFunctionPrologue(TargetCodeRegion *TCR,
         }
       }
     }
-    Out << (*i)->getDeclName().getAsString();
+    Out << VarName;
   }
   Out << ")\n{\n";
 
@@ -253,7 +259,10 @@ std::string TargetCode::generateFunctionName(TargetCodeRegion *TCR) {
 
 void TargetCode::handleArrays(const clang::ArrayType **t,
                               std::list<std::string> &DimString, int &dim,
-                              TargetCodeRegion *TCR, std::string &elemType) {
+                              std::vector<int> &VariableDims,
+                              TargetCodeRegion *TCR,
+                              std::string &elemType,
+                              const std::string &ArrayName) {
   auto OrigT = *t;
 
   if (!t) {
@@ -274,9 +283,9 @@ void TargetCode::handleArrays(const clang::ArrayType **t,
     DEBUGP("ArrayType VAT");
     std::string PrettyStr = "";
     llvm::raw_string_ostream PrettyOS(PrettyStr);
-    clang::PrintingPolicy PP(TCR->GetLangOpts());
-    t1->getSizeExpr()->printPretty(PrettyOS, NULL, PP);
+    PrettyOS << "__sotoc_vla_dim" << dim << "_" << ArrayName;
     DimString.push_back(PrettyOS.str());
+    VariableDims.push_back(dim);
     ++dim;
 
   } else {
@@ -290,6 +299,6 @@ void TargetCode::handleArrays(const clang::ArrayType **t,
       OrigT->getElementType().getTypePtr());
   if (*t) {
     // Recursively handle all dimensions
-    handleArrays(t, DimString, dim, TCR, elemType);
+    handleArrays(t, DimString, dim, VariableDims, TCR, elemType, ArrayName);
   }
 }
diff --git a/clang/tools/sotoc/src/TargetCode.h b/clang/tools/sotoc/src/TargetCode.h
@@ -85,5 +85,7 @@ 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,
-                    TargetCodeRegion *TCR, std::string &elemType);
+                    std::vector<int> &VariableDimensions,
+                    TargetCodeRegion *TCR, std::string &elemType,
+                    const std::string &ArrayName);
 };
diff --git a/clang/tools/sotoc/test/CMakeLists.txt b/clang/tools/sotoc/test/CMakeLists.txt
@@ -33,7 +33,9 @@ if(NOT SOTOC_LLVM_LIT_EXECUTABLE)
 endif()
 
 
-set(SOTOC_LIT_ARGS -v --no-progress-bar)
+set(SOTOC_LIT_ARGS "-v;--no-progress-bar;-j 4")
+
+set(SOTOC_FILECHECK_EXE ${LLVM_FILECHECK_EXE})
 
 add_custom_target(check-sotoc
   COMMAND ${PYTHON_EXECUTABLE} ${SOTOC_LLVM_LIT_EXECUTABLE} ${SOTOC_LIT_ARGS} ${CMAKE_CURRENT_BINARY_DIR}

diff --git a/clang/tools/sotoc/test/arrays/a.out b/clang/tools/sotoc/test/arrays/a.out
diff --git a/clang/tools/sotoc/test/arrays/global_array_static.c b/clang/tools/sotoc/test/arrays/global_array_static.c
@@ -12,12 +12,14 @@
 
  int tmp = 0;
 
- #pragma omp target device(0) map(from:X[:10])
+ #pragma omp target update to(X[:10])
+ #pragma omp target device(0)
  {
   for(int i = 0; i < 10; ++i){
    X[i] = 1;
   }
  }
+ #pragma omp target update from(X[:10])
 
  for(int i = 0; i < 10; ++i){
    tmp += X[i];

diff --git a/clang/tools/sotoc/test/arrays/static_variable_length_array_2d.c b/clang/tools/sotoc/test/arrays/static_variable_length_array_2d.c
@@ -7,7 +7,7 @@ int main(){
   int size=512;
   float A[size][2];
 
-  #pragma omp target map(tofrom:A[:size*2])
+  #pragma omp target map(tofrom:A[:size][:2])
   {
     int i;
     for(i=0; i< size; i++){
@@ -17,7 +17,11 @@ int main(){
   }
 
   for (j = 0; j < size; j+=64) {
+<<<<<<< HEAD
+    printf("%.2f %.2f ",A[j][0],A[j][1]);
+=======
     printf("%.2f %.2f",A[j][0],A[j][1]);
+>>>>>>> aurora_offloading_prototype
   }
   return 0;
 }

diff --git a/clang/tools/sotoc/test/arrays/static_variable_length_array_2d_sw.c b/clang/tools/sotoc/test/arrays/static_variable_length_array_2d_sw.c
@@ -0,0 +1,25 @@
+// RUN: %sotoc-transform-compile
+// RUN: %run-on-host | %filecheck %s
+
+#include <stdio.h>
+int main(){
+  int j;
+  int size=512;
+  float A[2][size];
+
+  #pragma omp target map(tofrom:A[:2][:size])
+  {
+    int i;
+    for(i=0; i< size; i++){
+      A[0][i]=i;
+      A[1][i]=i+1;
+    }
+  }
+
+  for (j = 0; j < size; j+=64) {
+    printf("%.2f %.2f ",A[0][j],A[1][j]);
+  }
+  return 0;
+}
+
+// CHECK: 0.00 1.00 64.00 65.00 128.00 129.00 192.00 193.00 256.00 257.00 320.00 321.00 384.00 385.00 448.00 449.00
diff --git a/clang/tools/sotoc/test/arrays/static_variable_length_array_3d.c b/clang/tools/sotoc/test/arrays/static_variable_length_array_3d.c
@@ -0,0 +1,28 @@
+// RUN: %sotoc-transform-compile
+// RUN: %run-on-host | %filecheck %s
+
+#include <stdio.h>
+int main(){
+  int j;
+  int size=512;
+  float A[size][2][5];
+
+  #pragma omp target map(tofrom:A[:size][:2][:5])
+  {
+    int k;
+    int i;
+    for(i=0; i< size; i++){
+      for(k=0; k< 5; k++){
+        A[i][0][k]=i;
+        A[i][1][k]=i+1;
+      }
+    }
+  }
+
+  for (j = 0; j < size; j+=64) {
+    printf("%.2f %.2f ",A[j][0][3],A[j][1][2]);
+  }
+  return 0;
+}
+
+// CHECK: 0.00 1.00 64.00 65.00 128.00 129.00 192.00 193.00 256.00 257.00 320.00 321.00 384.00 385.00 448.00 449.00
diff --git a/clang/tools/sotoc/test/arrays/static_variable_length_array_3d_sw.c b/clang/tools/sotoc/test/arrays/static_variable_length_array_3d_sw.c
@@ -0,0 +1,28 @@
+// RUN: %sotoc-transform-compile
+// RUN: %run-on-host | %filecheck %s
+
+#include <stdio.h>
+int main(){
+  int j;
+  int size=512;
+  float A[2][size][5];
+
+  #pragma omp target map(tofrom:A[:2][:size][:5])
+  {
+    int i;
+    int k;
+    for(i=0; i< size; i++){
+      for(k=0; k< 5; k++){
+        A[0][i][k]=i;
+        A[1][i][k]=i+1;
+      }
+    }
+  }
+
+  for (j = 0; j < size; j+=64) {
+    printf("%.2f %.2f ",A[0][j][3],A[1][j][2]);
+  }
+  return 0;
+}
+
+// CHECK: 0.00 1.00 64.00 65.00 128.00 129.00 192.00 193.00 256.00 257.00 320.00 321.00 384.00 385.00 448.00 449.00
diff --git a/clang/tools/sotoc/test/arrays/variable_length_array_1d.c b/clang/tools/sotoc/test/arrays/variable_length_array_1d.c
@@ -6,7 +6,7 @@ int main(){
   int j;
   int size=512;
   float A[size]; 
-  #pragma omp target map(tofrom:A[0:size],size)
+  #pragma omp target map(tofrom:A[0:size])
   {
     int i;
     for(i=0; i< size; i++){

diff --git a/clang/tools/sotoc/test/arrays/variable_length_array_2d.c b/clang/tools/sotoc/test/arrays/variable_length_array_2d.c
@@ -7,7 +7,8 @@ int main(){
   int sizeX=512;
   int sizeY=512;
   float A[sizeX][sizeY]; 
-  #pragma omp target map(tofrom:A[0:sizeX*sizeY])
+
+  #pragma omp target map(tofrom:A[:sizeX][:sizeY])
   {
     int i;
     int j;

diff --git a/clang/tools/sotoc/test/lit.site.cfg.in b/clang/tools/sotoc/test/lit.site.cfg.in
@@ -6,7 +6,7 @@ config.sotoc_obj_root = '@SOTOC_OBJ_ROOT@'
 config.library_dir = "@SOTOC_LIBOMP_LIBRARY_DIR@"
 config.omp_header_directory = "@SOTOC_OMP_HEADER_DIR@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
-config.filecheck = "@LLVM_FILECHECK_EXE@"
+config.filecheck = "@SOTOC_FILECHECK_EXE@"
 
 # Let the main config do the real work.
 lit_config.load_config(config, '@SOTOC_TEST_DIR@/lit.cfg')