From 47193fc3ee19bc44c290d87992229c0612ce51f6 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian.brown@amd.com>
Date: Mon, 17 Jun 2024 11:19:47 -0500
Subject: [PATCH] Split imperfect-loop-collapse into usm/non-usm variants

This patch splits the smoke-fails/imperfect-loop-collapse test into two
variants, one of which does and one which doesn't need unified shared
memory (which is orthogonal to the main purpose of the test).  It also
fixes the HSA_XNACK setting for smoke-fails/imperfect-loop-collapse-usm
version, so both versions now work.
---
 .../imperfect-loop-collapse-usm/Makefile      | 25 +++++++++++++
 .../imperfect_loop_collapse_usm.cpp           | 35 +++++++++++++++++++
 .../imperfect-loop-collapse/Makefile          |  5 +--
 .../imperfect_loop_collapse.cpp               |  5 ++-
 4 files changed, 63 insertions(+), 7 deletions(-)
 create mode 100644 test/smoke-fails/imperfect-loop-collapse-usm/Makefile
 create mode 100644 test/smoke-fails/imperfect-loop-collapse-usm/imperfect_loop_collapse_usm.cpp

diff --git a/test/smoke-fails/imperfect-loop-collapse-usm/Makefile b/test/smoke-fails/imperfect-loop-collapse-usm/Makefile
new file mode 100644
index 000000000..7b09b6a44
--- /dev/null
+++ b/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)
diff --git a/test/smoke-fails/imperfect-loop-collapse-usm/imperfect_loop_collapse_usm.cpp b/test/smoke-fails/imperfect-loop-collapse-usm/imperfect_loop_collapse_usm.cpp
new file mode 100644
index 000000000..4be83e0a5
--- /dev/null
+++ b/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;
+}
diff --git a/test/smoke-fails/imperfect-loop-collapse/Makefile b/test/smoke-fails/imperfect-loop-collapse/Makefile
index eb0665a22..7edc5fa32 100644
--- a/test/smoke-fails/imperfect-loop-collapse/Makefile
+++ b/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)
diff --git a/test/smoke-fails/imperfect-loop-collapse/imperfect_loop_collapse.cpp b/test/smoke-fails/imperfect-loop-collapse/imperfect_loop_collapse.cpp
index 4be83e0a5..11818bb4b 100644
--- a/test/smoke-fails/imperfect-loop-collapse/imperfect_loop_collapse.cpp
+++ b/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++)