[Openmp-commits] [openmp] 6368455 - [OpenMP][NFC] Add offloading tests for the new ompx APIs

[Johannes Doerfert via Openmp-commits]

Author: Johannes Doerfert
Date: 2023-07-31T15:45:53-07:00
New Revision: 63684550c4e25a1ba93b74179ca18b70c8696f1e

URL: https://github.com/llvm/llvm-project/commit/63684550c4e25a1ba93b74179ca18b70c8696f1e
DIFF: https://github.com/llvm/llvm-project/commit/63684550c4e25a1ba93b74179ca18b70c8696f1e.diff

LOG: [OpenMP][NFC] Add offloading tests for the new ompx APIs

Added: 
    openmp/libomptarget/test/offloading/ompx_coords.c
    openmp/libomptarget/test/offloading/ompx_saxpy_mixed.c

Modified: 
    

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/test/offloading/ompx_coords.c b/openmp/libomptarget/test/offloading/ompx_coords.c
new file mode 100644
index 00000000000000..46cb8d4e3c1c7c
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/ompx_coords.c
@@ -0,0 +1,61 @@
+// RUN: %libomptarget-compileopt-run-and-check-generic
+//
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+#include <omp.h>
+#include <ompx.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+struct info {
+  int tid, bid, tdim;
+};
+
+int main(int argc, char **argv) {
+  int N = 1 << 20;
+  if (argc > 1)
+    N = atoi(argv[1]);
+
+  struct info *X = (struct info *)malloc(sizeof(*X) * N);
+  memset(X, '0', sizeof(*X) * N);
+
+  int TL = 256;
+  int NT = (N + TL - 1) / TL;
+
+#pragma omp target data map(tofrom : X [0:N])
+#pragma omp target teams num_teams(NT) thread_limit(TL)
+  {
+#pragma omp parallel
+    {
+      int tid = ompx_thread_id_x();
+      int bid = ompx_block_id_x();
+      int tdim = ompx_thread_dim_x();
+      int gid = tid + bid * tdim;
+      if (gid < N) {
+        X[gid].tid = tid;
+        X[gid].bid = bid;
+        X[gid].tdim = tdim;
+      };
+    }
+  }
+
+  int tid = 0, bid = 0, tdim = 256;
+  for (int i = 0; i < N; i++) {
+    if (X[i].tid != tid || X[i].bid != bid || X[i].tdim != tdim) {
+      printf("%i: %i vs %i, %i vs %i, %i vs %i\n", i, X[i].tid, tid, X[i].bid,
+             bid, X[i].tdim, tdim);
+      return 1;
+    }
+    tid++;
+    if (tid == tdim) {
+      tid = 0;
+      bid++;
+    }
+  }
+
+  // CHECK: OK
+  printf("OK");
+  return 0;
+}

diff  --git a/openmp/libomptarget/test/offloading/ompx_saxpy_mixed.c b/openmp/libomptarget/test/offloading/ompx_saxpy_mixed.c
new file mode 100644
index 00000000000000..19c50f085a4afa
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/ompx_saxpy_mixed.c
@@ -0,0 +1,55 @@
+// RUN: %libomptarget-compileopt-run-and-check-generic
+//
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+#include <math.h>
+#include <omp.h>
+#include <ompx.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main(int argc, char **argv) {
+  int N = 1 << 29;
+  if (argc > 1)
+    N = atoi(argv[1]);
+  float a = 2.f;
+
+  float *X = (float *)malloc(sizeof(*X) * N);
+  float *Y = (float *)malloc(sizeof(*X) * N);
+
+  for (int i = 0; i < N; i++) {
+    X[i] = 1.0f;
+    Y[i] = 2.0f;
+  }
+
+  int TL = 256;
+  int NT = (N + TL - 1) / TL;
+
+#pragma omp target data map(to : X [0:N]) map(Y [0:N])
+#pragma omp target teams num_teams(NT) thread_limit(TL)
+  {
+#pragma omp parallel
+    {
+      int tid = ompx_thread_id_x();
+      int bid = ompx_block_id_x();
+      int tdim = ompx_thread_dim_x();
+      int gid = tid + bid * tdim;
+      if (gid < N)
+        Y[gid] = a * X[gid] + Y[gid];
+    }
+  }
+
+  float maxError = 0.0f;
+  for (int i = 0; i < N; i++) {
+    maxError = fmax(maxError, fabs(Y[i] - 4.0f));
+    if (maxError) {
+      printf("%i %f %f\n", i, maxError, Y[i]);
+      break;
+    }
+  }
+  // CHECK: Max error: 0.00
+  printf("Max error: %f\n", maxError);
+
+  return 0;
+}