Patches relating to OpenMP imperfect loop collapse

bin/rocm-test/scripts/parse_LLNL.sh

@@ -16,8 +16,14 @@ testname_regex='.*(test_\S*)'
 compile_regex='Compilation.*failed'
 runtime_regex='Running.+\.\.\.\s+([A-Z]*[a-z]*)'
 
-cd "$aompdir/test/LLNL/openmp5.0-tests"
-infile=`ls | grep "LLNL.run.log"`
+cd "$aompdir/test/LLNL/openmp5.0-tests" || exit 1
+declare -a infiles
+infiles=( LLNL.run.log* )
+if [ "${#infiles[@]}" -ne 1 ]; then
+  echo "Expected to find a single result file, bailing out" >&2
+  exit 1
+fi
+infile=${infiles[0]}
 
 # Clean up before parsing
 if [ -e passing-tests.txt ]; then
@@ -26,6 +32,9 @@ fi
 if [ -e failing-tests.txt ]; then
   rm failing-tests.txt
 fi
+if [ -e xfail-tests.txt ]; then
+  rm xfail-tests.txt
+fi
 if [ -e make-fail.txt ]; then
   rm make-fail.txt
 fi
@@ -49,6 +58,10 @@ function parse(){
       fi
       [[ "$line" =~ $testname_regex ]]
       llnltest=${BASH_REMATCH[1]}
+      if [ "$outfile" = "failing-tests.txt" ] && \
+         grep -Fq "$llnltest" xfail-list; then
+        outfile="xfail-tests.txt"
+      fi
       echo "$llnltest" >> $outfile
     fi
   done < "$infile"

test/LLNL/openmp5.0-tests/check_LLNL.sh

@@ -50,7 +50,15 @@ if [ "$1" == "log" ]; then
     log="LLNL.run.log.$date"
   fi
   echo "Log enabled: $log"
- timeout 120 ./test.py $AOMP $AOMP_GPU 2>&1 | tee $log
+ timeout 480 ./test.py $AOMP $AOMP_GPU 2>&1 | tee $log
+ if [ ${PIPESTATUS[0]} -eq 124 ]; then
+   echo "WARNING: Testing timed out!" >&2
+   exit 1
+ fi
 else
- timeout 120 ./test.py $AOMP $AOMP_GPU
+ timeout 480 ./test.py $AOMP $AOMP_GPU
+ if [ $? -eq 124 ]; then
+   echo "WARNING: Testing timed out!" >&2
+   exit 1
+ fi
 fi

test/LLNL/openmp5.0-tests/test.py

@@ -19,13 +19,19 @@
 
 def get_tests(file_name):
     d=defaultdict(list)
+    o=defaultdict(lambda: '')
     with open(file_name,"r") as infile:
         for line in infile.readlines():
             #print(line)
-            sline=line.split()
+            if '/' in line:
+                files,opts = line.split('/')
+                sline=files.split()
+                o[sline[0][:-4]]=opts
+            else:
+                sline=line.split()
             for s in sline:
                 d[sline[0][:-4]].append(s)
-    return d
+    return d,o
 
 def run(tests):
     pass_count=0
@@ -35,40 +41,49 @@ def run(tests):
         cmd="./"+t
         cmd_s=cmd.split()
         p4=subprocess.Popen(cmd_s,stdout=subprocess.PIPE,stderr=subprocess.STDOUT)
-        for lines in p4.stdout.readlines():
+        stdout,stderr = p4.communicate()
+        stdout = stdout.decode('utf-8')
+        for lines in stdout.splitlines():
             #print(lines)
-            if b"error" in lines:
+            if "error" in lines:
                 print("Failed")
                 passs=False
                 break
-            if b"FAIL" in lines:
+            if "FAIL" in lines:
                 print("Failed")
                 passs=False
                 break
+        if p4.returncode != 0:
+            passs=False
         if passs: 
             print ("Passed")
             pass_count+=1
     return pass_count
 
-def compile(CC,LIBS, tests):
+def compile(CC, LIBS, tests, opts):
 
     runnables=[]
     for key,value in tests.items():
         with open(key+".ERR","w") as efile:
             passs=True
             for v in value:
                 fail=False
-                cmd=CC+" -c "+ v
+                extraopts=opts[key]
+                cmd=CC+" -c " + v + " " + extraopts
                 cmd_s=cmd.split()
                 p2=subprocess.Popen(cmd_s,stdout=subprocess.PIPE,stderr=subprocess.STDOUT)
-                for lines in p2.stdout.readlines():
-                    efile.write(lines.decode('utf8'))
-                    if b"error" in lines:
+                stdout,stderr = p2.communicate()
+                stdout = stdout.decode('utf-8')
+                for lines in stdout.splitlines():
+                    efile.write(lines)
+                    if "error" in lines:
                         fail=True
                         print("Compilation of ",v," failed")
                         passs = passs and not fail
                         break
                     passs = passs and not fail
+                if p2.returncode != 0:
+                    passs=False
 
             if passs:
                 print("Compiling ",key)
@@ -79,20 +94,23 @@ def compile(CC,LIBS, tests):
                 #print("Final compile command is ",cmd)
                 cmd_s=cmd.split()
                 p3=subprocess.Popen(cmd_s,stdout=subprocess.PIPE,stderr=subprocess.STDOUT)
-                for lines in p3.stdout.readlines():
-                    print(lines.decode('utf8'))
-                    if b"error" in lines:
+                stdout, stderr = p3.communicate()
+                stdout = stdout.decode('utf-8')
+                for lines in stdout.splitlines():
+                    print(lines)
+                    if "error" in lines:
                         print("Linking of ",v," failed\n")
                         break
-                runnables.append(key)
+                if p3.returncode == 0:
+                    runnables.append(key)
     return runnables
 def main():
-    tests=get_tests("test_list")
+    tests,opts=get_tests("test_list")
 # Change Compile line in CC and LIBS
     CC="{}/bin/clang++  -O2  -target x86_64-pc-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march={}".format(AOMP, AOMP_GPU)
     LIBS = ""
 # End Compile line 
-    runnables=compile(CC,LIBS, tests)
+    runnables=compile(CC, LIBS, tests, opts)
     print("\nRunnable tests are:")
     for r in runnables:
         print(r)

test/LLNL/openmp5.0-tests/test_imperfect_loop_warning.cxx

@@ -0,0 +1,34 @@
+/* Scan error output from the compilation of "test_imperfect_loop.cxx".
+   Every problem is a hammer here.  */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+int main (void) {
+  FILE *f = fopen ("test_imperfect_loop.ERR", "r");
+  size_t bufsize = 1024;
+  char *buf = (char*) malloc (bufsize);
+  bool found = false;
+
+  if (!f) {
+    perror ("fopen");
+    goto fail;
+  }
+  while (!feof (f) && !found) {
+    getline (&buf, &bufsize, f);
+    if (strstr (buf, "remark: Collapsing imperfectly-nested loop may "
+                     "introduce unexpected data dependencies"))
+      found = true;
+  }
+  fclose (f);
+
+  if (found) {
+    printf ("Imperfect loop warning test:: Pass\n");
+    exit (0);
+  }
+
+fail:
+  printf ("Imperfect loop warning test:: FAIL\n");
+  return 1;
+}

test/LLNL/openmp5.0-tests/test_list

@@ -3,7 +3,8 @@ test_allocate_directive.cxx
 test_concurrent_map.cxx
 test_implicit_declare_target.cxx foo.cxx
 test_declare_variant.cxx
-test_imperfect_loop.cxx
+test_imperfect_loop.cxx / -Rpass-analysis=openmp-opt
+test_imperfect_loop_warning.cxx
 test_loop.cxx
 test_lvalue_map_motion.cxx
 test_non_rectangular.cxx

test/LLNL/openmp5.0-tests/xfail-list

@@ -0,0 +1 @@
+test_imperfect_loop run

test/smoke-fails/imperfect-loop-collapse-usm/Makefile

@@ -0,0 +1,25 @@
+include ../../Makefile.defs
+
+TESTNAME     = imperfect_loop_collapse_usm
+TESTSRC_MAIN = imperfect_loop_collapse_usm.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+HSA_XNACK ?= 1
+SUPPORTED    = $(SUPPORTS_USM)
+
+# Our "run" target gets overridden.  Make sure we run with HSA_XNACK set
+# appropriately.
+RUNENV	     += HSA_XNACK=${HSA_XNACK}
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
+
+run:
+	HSA_XNACK=${HSA_XNACK} ./$(TESTNAME)

test/smoke-fails/imperfect-loop-collapse-usm/imperfect_loop_collapse_usm.cpp

@@ -0,0 +1,35 @@
+#include<cstdio>
+
+#define N 1024
+
+#pragma omp requires unified_shared_memory
+
+int main() {
+  double *a = new double[N*N];
+
+  #pragma omp parallel for collapse(2)
+  for(int i = 0; i < N; i++)
+    for(int j = 0; j < N; j++)
+      a[i*N+j] = (double) (i*N+j);
+
+  #pragma omp target teams distribute parallel for collapse(2)
+  for(int i = 0; i < N; i++) {
+    double k = i*3.14;
+    for(int j = 0; j < N; j++)
+      a[i*N+j] += k;
+  }
+
+  //check
+  int err = 0;
+  for(int i = 0; i < N; i++) {
+    double k = i*3.14;
+    for(int j = 0; j < N; j++)
+      if (a[i*N+j] != (double) (i*N+j) + k) {
+        err++;
+        printf("Error at (%d,%d): got %lf expected %lf\n", i, j, a[i*N+j], (double) (i*N+j) + k);
+        if (err > 10) return err;
+      }
+  }
+
+  return err;
+}

test/smoke-fails/imperfect-loop-collapse/Makefile

@@ -9,13 +9,10 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-HSA_XNACK ?= 1
-SUPPORTED    = $(SUPPORTS_USM)
-
 #-ccc-print-phases
 #"-\#\#\#"
 
 include ../Makefile.rules
 
 run:
-	HSA_XNACK=${HSA_XNACK} ./$(TESTNAME)
+	./$(TESTNAME)

test/smoke-fails/imperfect-loop-collapse/imperfect_loop_collapse.cpp

@@ -2,8 +2,6 @@
 
 #define N 1024
 
-#pragma omp requires unified_shared_memory
-
 int main() {
   double *a = new double[N*N];
 
@@ -12,7 +10,8 @@ int main() {
     for(int j = 0; j < N; j++)
       a[i*N+j] = (double) (i*N+j);
 
-  #pragma omp target teams distribute parallel for collapse(2)
+  #pragma omp target teams distribute parallel for collapse(2) \
+                     map(tofrom: a[0:N*N])
   for(int i = 0; i < N; i++) {
     double k = i*3.14;
     for(int j = 0; j < N; j++)