Port SYCL tests to OpenMP for AddressSanitizer

offload/test/sanitizer/sycl-tests/bad-free/bad-free-host.c

@@ -0,0 +1,13 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int FakePtr[3] = {1, 2, 3};
+  int Device = omp_get_default_device();
+  omp_target_free(FakePtr, Device);
+  return 0;
+}

offload/test/sanitizer/sycl-tests/bad-free/bad-free-minus1.c

@@ -0,0 +1,18 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 100;
+  int N_SZ = sizeof(int) * N;
+
+  int Device = omp_get_default_device();
+
+  int *DevPtr = (int *)omp_target_alloc(N_SZ, Device);
+  omp_target_free(DevPtr - 1, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/bad-free/bad-free-plus1.c

@@ -0,0 +1,18 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 100;
+  int N_SZ = sizeof(int) * N;
+
+  int Device = omp_get_default_device();
+
+  int *DevPtr = (int *)omp_target_alloc(N_SZ, Device);
+  omp_target_free(DevPtr + 1, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/double-free/double-free.c

@@ -0,0 +1,19 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 100;
+  int N_SZ = sizeof(int) * N;
+
+  int Device = omp_get_default_device();
+
+  int *DevPtr = (int *)omp_target_alloc(N_SZ, Device);
+  omp_target_free(DevPtr, Device);
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/DeviceGlobal/device-global.c

@@ -0,0 +1,20 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+// Port of https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/
+// AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp
+
+#include <omp.h>
+#include <stdio.h>
+
+#define ITEM_COUNT 3
+
+char dev_global[5];
+#pragma omp declare target(dev_global)
+
+int main() {
+#pragma omp target
+  { dev_global[8] = 42; }
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/USM/parallel-for-char.c

@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 10;
+  int N_SZ = sizeof(char) * N;
+
+  int Device = omp_get_default_device();
+
+  char *DevPtr = (char *)omp_target_alloc(N_SZ, Device);
+
+#pragma omp target is_device_ptr(DevPtr)
+  {
+#pragma omp parallel for
+    for (int i = 0; i < N; i++) {
+      DevPtr[i] = '*';
+    }
+  }
+
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/USM/parallel-for-double.c

@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 10;
+  int N_SZ = sizeof(double) * N;
+
+  int Device = omp_get_default_device();
+
+  double *DevPtr = (double *)omp_target_alloc(N_SZ, Device);
+
+#pragma omp target is_device_ptr(DevPtr)
+  {
+#pragma omp parallel for
+    for (int i = 0; i < N; i++) {
+      DevPtr[i] = 1.23;
+    }
+  }
+
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/USM/parallel-for-func.c

@@ -0,0 +1,29 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+__attribute__((noinline)) void foo(int *array, int i) { array[i] = 1; }
+
+int main() {
+  int N = 10;
+  int N_SZ = sizeof(int) * N;
+
+  int Device = omp_get_default_device();
+
+  int *DevPtr = (int *)omp_target_alloc(N_SZ, Device);
+
+#pragma omp target is_device_ptr(DevPtr)
+  {
+#pragma omp parallel for
+    for (int i = 0; i < N; i++) {
+      foo(DevPtr, i);
+    }
+  }
+
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/USM/parallel-for-int.c

@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 10;
+  int N_SZ = sizeof(int) * N;
+
+  int Device = omp_get_default_device();
+
+  int *DevPtr = (int *)omp_target_alloc(N_SZ, Device);
+
+#pragma omp target is_device_ptr(DevPtr)
+  {
+#pragma omp parallel for
+    for (int i = 0; i < N; i++) {
+      DevPtr[i] = 2;
+    }
+  }
+
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/USM/parallel-for-short.c

@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = 10;
+  int N_SZ = sizeof(short) * N;
+
+  int Device = omp_get_default_device();
+
+  short *DevPtr = (short *)omp_target_alloc(N_SZ, Device);
+
+#pragma omp target is_device_ptr(DevPtr)
+  {
+#pragma omp parallel for
+    for (int i = 0; i < N; i++) {
+      DevPtr[i] = 2;
+    }
+  }
+
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}

offload/test/sanitizer/sycl-tests/out-of-bounds/buffer/buffer1d.c

@@ -0,0 +1,30 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define X 3
+
+int main() {
+  int A[X] = {1, 2, 3};
+
+#pragma omp target map(tofrom : A)
+  {
+    for (int i = 0; i < X; i++) {
+      A[i] = A[i] * 2;
+    }
+  }
+
+  for (int i = 0; i < X; i++) {
+    printf("%d\n", A[i]);
+  }
+
+  return 0;
+}
+
+// CHECK: 2
+// CHECK: 4
+// CHECK: 6

offload/test/sanitizer/sycl-tests/out-of-bounds/buffer/buffer2d.c

@@ -0,0 +1,38 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define X 2
+#define Y 3
+
+int main() {
+  int A[X][Y] = {{1, 2, 3}, {4, 5, 6}};
+
+#pragma omp target map(tofrom : A)
+  {
+    for (int i = 0; i < X; i++) {
+      for (int j = 0; j < Y; j++) {
+        A[i][j] = A[i][j] * 2;
+      }
+    }
+  }
+
+  for (int i = 0; i < X; i++) {
+    for (int j = 0; j < Y; j++) {
+      printf("%d", A[i][j]);
+      if (j + 1 != Y) {
+        printf(" ");
+      }
+    }
+    printf("\n");
+  }
+
+  return 0;
+}
+
+// CHECK: 2 4 6
+// CHECK-NEXT: 8 10 12

offload/test/sanitizer/sycl-tests/out-of-bounds/buffer/buffer3d.c

@@ -0,0 +1,55 @@
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload
+// RUN: %libomptarget-run-generic 2>&1 > %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define X 4
+#define Y 3
+#define Z 2
+
+int main() {
+  int A[X][Y][Z] = {
+      {{0, 1}, {2, 3}, {4, 5}},
+      {{6, 7}, {8, 9}, {10, 11}},
+      {{12, 13}, {14, 15}, {16, 17}},
+      {{18, 19}, {20, 21}, {22, 23}},
+  };
+
+#pragma omp target map(tofrom : A)
+  {
+    for (int i = 0; i < X; i++) {
+      for (int j = 0; j < Y; j++) {
+        for (int k = 0; k < Z; k++) {
+          A[i][j][k] = A[i][j][k] * 2;
+        }
+      }
+    }
+  }
+
+  for (int i = 0; i < X; i++) {
+    for (int j = 0; j < Y; j++) {
+      printf("(");
+      for (int k = 0; k < Z; k++) {
+        printf("%d", A[i][j][k]);
+        if (k + 1 != Z) {
+          printf(",");
+        }
+      }
+      printf(")");
+      if (j + 1 != Y) {
+        printf(",");
+      }
+    }
+    printf("\n");
+  }
+
+  return 0;
+}
+
+// CHECK: (0,2),(4,6),(8,10)
+// CHECK: (12,14),(16,18),(20,22)
+// CHECK: (24,26),(28,30),(32,34)
+// CHECK: (36,38),(40,42),(44,46)