From 6a25bf3de6a567adaa58e901a99d5b5f35b87c35 Mon Sep 17 00:00:00 2001
From: Amit Pandey <amit.pandey@amd.com>
Date: Wed, 10 Apr 2024 13:52:20 +0000
Subject: [PATCH 1/6] Add Smoke tests for ASan.

---
 test/Makefile.defs                            | 11 ++-
 test/smoke-asan/Makefile                      | 89 +++++++++++++++++++
 test/smoke-asan/Makefile.rules                |  1 +
 test/smoke-asan/check_smoke_asan.sh           |  1 +
 .../hip-global-buffer-overflow/Makefile       | 30 +++++++
 .../hip-global-buffer-overflow.cpp            | 37 ++++++++
 .../hip-heap-buffer-overflow/Makefile         | 30 +++++++
 .../hip-heap-buffer-overflow.cpp              | 40 +++++++++
 test/smoke-asan/hip-use-after-free/Makefile   | 30 +++++++
 .../hip-use-after-free/hip-use-after-free.cpp | 40 +++++++++
 .../omp-global-buffer-overflow/Makefile       | 21 +++++
 .../omp-global-buffer-overflow.cpp            | 22 +++++
 .../omp-heap-buffer-overflow/Makefile         | 21 +++++
 .../omp-heap-buffer-overflow.cpp              | 18 ++++
 14 files changed, 389 insertions(+), 2 deletions(-)
 create mode 100644 test/smoke-asan/Makefile
 create mode 120000 test/smoke-asan/Makefile.rules
 create mode 120000 test/smoke-asan/check_smoke_asan.sh
 create mode 100644 test/smoke-asan/hip-global-buffer-overflow/Makefile
 create mode 100644 test/smoke-asan/hip-global-buffer-overflow/hip-global-buffer-overflow.cpp
 create mode 100644 test/smoke-asan/hip-heap-buffer-overflow/Makefile
 create mode 100644 test/smoke-asan/hip-heap-buffer-overflow/hip-heap-buffer-overflow.cpp
 create mode 100644 test/smoke-asan/hip-use-after-free/Makefile
 create mode 100644 test/smoke-asan/hip-use-after-free/hip-use-after-free.cpp
 create mode 100644 test/smoke-asan/omp-global-buffer-overflow/Makefile
 create mode 100644 test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
 create mode 100644 test/smoke-asan/omp-heap-buffer-overflow/Makefile
 create mode 100644 test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp

diff --git a/test/Makefile.defs b/test/Makefile.defs
index e12278e9d..a7f7ae430 100644
--- a/test/Makefile.defs
+++ b/test/Makefile.defs
@@ -210,6 +210,7 @@ MPILIBS ?= -L${OMPIDIR}/lib$(MPI64) -lmpi
     endif
 
 AOMP_GPU       ?= $(INSTALLED_GPU)
+AOMP_SANITIZER ?= $(AOMP_BUILD_SANITIZER)
 CC              = $(AOMP)/bin/clang
 OG11           ?= $(HOME)/git/og11/install
 OG11FLAGS      ?= -O3 -fopenmp -foffload=-march=$(AOMP_GPU) -I. -lgfortran -L $(OG11)/lib64
@@ -237,6 +238,12 @@ ifeq (gfx11,$(findstring gfx11,$(AOMP_GPU)))
  AOMP_WSIZE = -DWAVE_SIZE=32
 endif
 
+ifeq ($(AOMP_SANITIZER),)
+ASAN_FLAGS = -fsanitize=address -shared-libasan -g
+# For Sanitizer below environment variable should always be enabled by default
+AOMP_TARGET_FEATURES=:xnack+
+endif
+
 ifeq ($(AOMP_TARGET_FEATURES),)
 GPU_W_FEATURES = $(AOMP_GPU)
 else
@@ -247,9 +254,9 @@ USE_OFFLOAD_ARCH ?= 1
 ifeq ($(TARGET),)
   ifeq ($(OMP_HOST),)
    ifeq ($(USE_OFFLOAD_ARCH),1)
-	TARGET  =  --offload-arch=$(GPU_W_FEATURES)
+	TARGET  =  --offload-arch=$(GPU_W_FEATURES) $(ASAN_FLAGS)
    else
-	TARGET  =  -fopenmp-targets=$(AOMP_GPUTARGET) -Xopenmp-target=$(AOMP_GPUTARGET) -march=$(GPU_W_FEATURES)
+	TARGET  =  -fopenmp-targets=$(AOMP_GPUTARGET) -Xopenmp-target=$(AOMP_GPUTARGET) -march=$(GPU_W_FEATURES) $(ASAN_FLAGS)
    endif
   endif
 endif
diff --git a/test/smoke-asan/Makefile b/test/smoke-asan/Makefile
new file mode 100644
index 000000000..3577ffb8b
--- /dev/null
+++ b/test/smoke-asan/Makefile
@@ -0,0 +1,89 @@
+include ../Makefile.defs
+
+TESTS_DIR = \
+    hip-global-buffer-overflow \
+    hip-heap-buffer-overflow \
+		hip-use-after-free \
+		omp-global-buffer-overflow \
+		omp-heap-buffer-overflow
+
+all:
+	@for test_dir in $(TESTS_DIR); do \
+	  echo; \
+	  test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
+	  echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir"; \
+	  $(MAKE) -C $$test_dir; \
+	done
+
+run run_obin run_sbin run_llbin clean clean_log llbin sbin obin:
+	@for test_dir in $(TESTS_DIR); do \
+	  echo $$nnn; \
+	  test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
+	  echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir $@"; \
+	  $(MAKE) -C $$test_dir $@; \
+	done
+
+check:
+	 @for test_dir in $(TESTS_DIR); do \
+          echo $$nnn; \
+          test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
+          echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir $@"; \
+          $(MAKE) -C $$test_dir $@; \
+        done
+
+.ll .ll.s .ll.o .s .s.o .o:
+	@for test_dir in $(TESTS_DIR); do \
+	  echo $$nnn; \
+	  test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
+	  echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir $$test_name$@"; \
+	  $(MAKE) -C $$test_dir $$test_name$@; \
+	done
+
+help:
+	@echo
+	@echo "LLVM Tool Chain:	$(AOMP)/bin"
+	@echo "Offload Targets:	$(TARGET)"
+	@echo "Host Target:		$(AOMP_CPUTARGET)"
+	@echo "Application Dirs:	$(TESTS_DIR)"
+	@echo
+	@echo "This Makefile supports the following flags:"
+	@echo
+	@echo "	make llbin      // Link pass only"
+	@echo "	make run_llbin  // Execute llbin"
+	@echo
+	@echo "	make sbin       // Link pass only"
+	@echo "	make run_sbin   // Execute sbin"
+	@echo
+	@echo "	make obin       // Link pass only"
+	@echo "	make run_obin   // Execute obin"
+	@echo
+	@echo "	make .ll        // Compile pass only             : -c -S -emit-llvm"
+	@echo "	make .ll.s      // Backend pass only             : -c -S"
+	@echo "	make .ll.o      // Assemble pass only            : -c"
+	@echo "	make .s         // Compile & Backend passes      : -c -S"
+	@echo "	make .s.o       // Assemble pass only            : -c"
+	@echo "	make .o         // Compile, Backend, Assemble    : -c"
+	@echo
+	@echo "	make            // All passes, build all examples from Application Dirs"
+	@echo "	make run        // Execute all binaries from Application Dirs"
+	@echo
+	@echo "	make clean"
+	@echo "	make clean_log"
+	@echo "	make help"
+	@echo
+	@echo " Environment variables to control compilation & execution"
+	@echo "    VERBOSE=1        See lots of compiler messages and driver actions"
+	@echo "    TEMPS=1          Do not delete intermediate files"
+	@echo "    OFFLOAD_DEBUG=1  See Runtime diagnostics for each call to libomptarget API"
+	@echo "    TARGET=          Override Makefile target"
+	@echo
+	@echo " Compile Environment:  $(SETENV)"
+	@echo
+	@echo " Run Environment:      $(RUNENV)"
+	@echo
+	@echo " Compile Flags:        $(CFLAGS) $(EXTRA_CFLAGS)"
+	@echo
+	@echo " OMP Compile Flags:    $(OMP_FLAGS) $(EXTRA_OMP_FLAGS)"
+	@echo
+	@echo " Link Flags:           $(LINK_FLAGS) $(EXTRA_LDFLAGS)"
+	@echo
diff --git a/test/smoke-asan/Makefile.rules b/test/smoke-asan/Makefile.rules
new file mode 120000
index 000000000..b96345353
--- /dev/null
+++ b/test/smoke-asan/Makefile.rules
@@ -0,0 +1 @@
+../smoke/Makefile.rules
\ No newline at end of file
diff --git a/test/smoke-asan/check_smoke_asan.sh b/test/smoke-asan/check_smoke_asan.sh
new file mode 120000
index 000000000..9beff5fdb
--- /dev/null
+++ b/test/smoke-asan/check_smoke_asan.sh
@@ -0,0 +1 @@
+../smoke/check_smoke.sh
\ No newline at end of file
diff --git a/test/smoke-asan/hip-global-buffer-overflow/Makefile b/test/smoke-asan/hip-global-buffer-overflow/Makefile
new file mode 100644
index 000000000..2c94934fc
--- /dev/null
+++ b/test/smoke-asan/hip-global-buffer-overflow/Makefile
@@ -0,0 +1,30 @@
+include ../../Makefile.defs
+
+TESTNAME     = hip-global-buffer-overflow
+TESTSRC_MAIN = hip-global-buffer-overflow.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+VERS = $(shell $(AOMP)/bin/clang --version | grep -oP '(?<=clang version )[0-9]+')
+ifeq ($(shell expr $(VERS) \>= 12.0), 1)
+  RPTH = -Wl,-rpath,$(AOMPHIP)/lib/asan
+  LLIB = -L$(AOMPHIP)/lib/asan
+endif
+
+CFLAGS = -x hip -std=c++11 $(LLIB) -lamdhip64 $(RPTH)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/hip-global-buffer-overflow/hip-global-buffer-overflow.cpp b/test/smoke-asan/hip-global-buffer-overflow/hip-global-buffer-overflow.cpp
new file mode 100644
index 000000000..dcc01939b
--- /dev/null
+++ b/test/smoke-asan/hip-global-buffer-overflow/hip-global-buffer-overflow.cpp
@@ -0,0 +1,37 @@
+#define __HIP_PLATFORM_AMD__ 1
+#include <hip/hip_runtime.h>
+
+#define N 100
+
+void printHipError(hipError_t error) {
+  printf("Hip Error: %s\n", hipGetErrorString(error));
+}
+
+bool hipCallSuccessfull(hipError_t error) {
+  if (error != hipSuccess)
+    printHipError(error);
+  return error == hipSuccess;
+}
+
+__device__ int D_Ptr[N];
+
+__global__ void Initialize(int n) {
+  int index = blockDim.x * blockIdx.x + threadIdx.x;
+  if (index < n) {
+    D_Ptr[index + 1] = 2 * (index + 1);
+  }
+}
+
+int main(int argc, char *argv[]) {
+  size_t NBytes = N * sizeof(int);
+  int NumOfThreadBlocks = (N + 64 - 1) / 64;
+  int ThreadBlockSize = 64;
+  hipLaunchKernelGGL(Initialize, dim3(NumOfThreadBlocks), dim3(ThreadBlockSize),
+                     0, 0, N);
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: global-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in Initialize(int) at {{.*}}aomp/test/smoke-asan/hip-global-buffer-overflow/hip-global-buffer-overflow.cpp:21:15
diff --git a/test/smoke-asan/hip-heap-buffer-overflow/Makefile b/test/smoke-asan/hip-heap-buffer-overflow/Makefile
new file mode 100644
index 000000000..6c1a9c4fd
--- /dev/null
+++ b/test/smoke-asan/hip-heap-buffer-overflow/Makefile
@@ -0,0 +1,30 @@
+include ../../Makefile.defs
+
+TESTNAME     = hip-heap-buffer-overflow
+TESTSRC_MAIN = hip-heap-buffer-overflow.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+VERS = $(shell $(AOMP)/bin/clang --version | grep -oP '(?<=clang version )[0-9]+')
+ifeq ($(shell expr $(VERS) \>= 12.0), 1)
+  RPTH = -Wl,-rpath,$(AOMPHIP)/lib/asan
+  LLIB = -L$(AOMPHIP)/lib/asan
+endif
+
+CFLAGS = -x hip -std=c++11 $(LLIB) -lamdhip64 $(RPTH)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/hip-heap-buffer-overflow/hip-heap-buffer-overflow.cpp b/test/smoke-asan/hip-heap-buffer-overflow/hip-heap-buffer-overflow.cpp
new file mode 100644
index 000000000..7d9d38e15
--- /dev/null
+++ b/test/smoke-asan/hip-heap-buffer-overflow/hip-heap-buffer-overflow.cpp
@@ -0,0 +1,40 @@
+#define __HIP_PLATFORM_AMD__ 1
+#include <hip/hip_runtime.h>
+
+void printHipError(hipError_t error) {
+  printf("Hip Error: %s\n", hipGetErrorString(error));
+}
+
+bool hipCallSuccessfull(hipError_t error) {
+  if (error != hipSuccess)
+    printHipError(error);
+  return error == hipSuccess;
+}
+
+__global__ void Initialize(int n, int *ptr) {
+  int index = blockDim.x * blockIdx.x + threadIdx.x;
+  if (index < n) {
+    ptr[index + 1] = 2 * (index + 1);
+  }
+}
+
+int main(int argc, char *argv[]) {
+  int N = 100;
+  size_t NBytes = N * sizeof(int);
+  int *H_Ptr = new int[N];
+  int *D_Ptr;
+  int NumOfThreadBlocks = (N + 64 - 1) / 64;
+  int ThreadBlockSize = 64;
+  hipCallSuccessfull(hipMalloc(&D_Ptr, NBytes));
+  hipLaunchKernelGGL(Initialize, dim3(NumOfThreadBlocks), dim3(ThreadBlockSize),
+                     0, 0, N, D_Ptr);
+  hipCallSuccessfull(hipMemcpy(H_Ptr, D_Ptr, NBytes, hipMemcpyDeviceToHost));
+  hipCallSuccessfull(hipFree(D_Ptr));
+  delete[] H_Ptr;
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: heap-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in Initialize(int, int*) at {{.*}}aomp/test/smoke-asan/hip-heap-buffer-overflow/hip-heap-buffer-overflow.cpp:17:13
diff --git a/test/smoke-asan/hip-use-after-free/Makefile b/test/smoke-asan/hip-use-after-free/Makefile
new file mode 100644
index 000000000..afd6bc684
--- /dev/null
+++ b/test/smoke-asan/hip-use-after-free/Makefile
@@ -0,0 +1,30 @@
+include ../../Makefile.defs
+
+TESTNAME     = hip-use-after-free
+TESTSRC_MAIN = hip-use-after-free.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+VERS = $(shell $(AOMP)/bin/clang --version | grep -oP '(?<=clang version )[0-9]+')
+ifeq ($(shell expr $(VERS) \>= 12.0), 1)
+  RPTH = -Wl,-rpath,$(AOMPHIP)/lib/asan
+  LLIB = -L$(AOMPHIP)/lib/asan
+endif
+
+CFLAGS = -x hip -std=c++11 $(LLIB) -lamdhip64 $(RPTH)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/hip-use-after-free/hip-use-after-free.cpp b/test/smoke-asan/hip-use-after-free/hip-use-after-free.cpp
new file mode 100644
index 000000000..ec8a50ee6
--- /dev/null
+++ b/test/smoke-asan/hip-use-after-free/hip-use-after-free.cpp
@@ -0,0 +1,40 @@
+#define __HIP_PLATFORM_AMD__ 1
+#include <hip/hip_runtime.h>
+
+void printHipError(hipError_t error) {
+  printf("Hip Error: %s\n", hipGetErrorString(error));
+}
+
+bool hipCallSuccessfull(hipError_t error) {
+  if (error != hipSuccess)
+    printHipError(error);
+  return error == hipSuccess;
+}
+
+__global__ void Initialize(int n, int *ptr) {
+  int index = blockDim.x * blockIdx.x + threadIdx.x;
+  if (index < n) {
+    ptr[index] = 2 * (index + 1);
+  }
+}
+
+int main(int argc, char *argv[]) {
+  int N = 100;
+  size_t NBytes = N * sizeof(int);
+  int *H_Ptr = new int[N];
+  int *D_Ptr;
+  int NumOfThreadBlocks = (N + 64 - 1) / 64;
+  int ThreadBlockSize = 64;
+  hipCallSuccessfull(hipMalloc(&D_Ptr, NBytes));
+  hipCallSuccessfull(hipFree(D_Ptr));
+  hipLaunchKernelGGL(Initialize, dim3(NumOfThreadBlocks), dim3(ThreadBlockSize),
+                     0, 0, N, D_Ptr);
+  hipCallSuccessfull(hipMemcpy(H_Ptr, D_Ptr, NBytes, hipMemcpyDeviceToHost));
+  delete[] H_Ptr;
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: heap-use-after-free on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in Initialize(int, int*) at {{.*}}aomp/test/smoke-asan/hip-use-after-free/hip-use-after-free.cpp:17:11
diff --git a/test/smoke-asan/omp-global-buffer-overflow/Makefile b/test/smoke-asan/omp-global-buffer-overflow/Makefile
new file mode 100644
index 000000000..6521eba03
--- /dev/null
+++ b/test/smoke-asan/omp-global-buffer-overflow/Makefile
@@ -0,0 +1,21 @@
+include ../../Makefile.defs
+
+TESTNAME     = omp-global-buffer-overflow
+TESTSRC_MAIN = omp-global-buffer-overflow.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp b/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
new file mode 100644
index 000000000..788f33701
--- /dev/null
+++ b/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
@@ -0,0 +1,22 @@
+#include <omp.h>
+#define N 100
+
+#pragma omp declare target
+int D_Ptr[N];
+#pragma omp end declare target
+
+int main(int argc, char *argv[]) {
+#pragma omp target data map(tofrom : D_Ptr[0 : N])
+  {
+#pragma omp target teams distribute parallel for
+    for (int i = 0; i < N; i++) {
+      D_Ptr[i + 1] = 2 * (i + 1);
+    }
+  }
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: global-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_fd00_68a0014_main_l11 at {{.*}}aomp/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp:13:13
diff --git a/test/smoke-asan/omp-heap-buffer-overflow/Makefile b/test/smoke-asan/omp-heap-buffer-overflow/Makefile
new file mode 100644
index 000000000..399089885
--- /dev/null
+++ b/test/smoke-asan/omp-heap-buffer-overflow/Makefile
@@ -0,0 +1,21 @@
+include ../../Makefile.defs
+
+TESTNAME     = omp-heap-buffer-overflow
+TESTSRC_MAIN = omp-heap-buffer-overflow.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp b/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp
new file mode 100644
index 000000000..eb54d8a84
--- /dev/null
+++ b/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp
@@ -0,0 +1,18 @@
+#include <omp.h>
+
+int main(int argc, char *argv[]) {
+  int N = 1000;
+  int *Ptr = new int[N];
+#pragma omp target data map(tofrom : Ptr[0 : N])
+#pragma omp target teams distribute parallel for
+  for (int i = 0; i < N; i++) {
+    Ptr[i + 1] = 2 * (i + 1);
+  }
+  delete[] Ptr;
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: heap-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_fd00_68a0017_main_l7 at {{.*}}aomp/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp:9:11

From bad8fcc139634847e7ec0cf10bda8eac0f63961f Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023@gmail.com>
Date: Tue, 16 Apr 2024 10:54:53 +0530
Subject: [PATCH 2/6] Address Comments.

1. Add a new makefile target 'check-asan'. 'check-asan' is capable of
   getting both the filecheck_status and test_status.

2. Remove dependency of AOMP_SANITIZER on AOMP_BUILD_SANITIZER.

3. Utilize the AOMP_SANITIZER environtment variable in check_smoke.sh to enable
   check-asan target for both parallel and sequential builds.
---
 test/Makefile.defs        |  9 ++++---
 test/smoke-asan/Makefile  |  6 ++---
 test/smoke/Makefile.rules | 49 ++++++++++++++++++++++++++++++++++++++-
 test/smoke/check_smoke.sh | 10 ++++++++
 4 files changed, 65 insertions(+), 9 deletions(-)

diff --git a/test/Makefile.defs b/test/Makefile.defs
index a7f7ae430..0b89e8f0d 100644
--- a/test/Makefile.defs
+++ b/test/Makefile.defs
@@ -210,7 +210,6 @@ MPILIBS ?= -L${OMPIDIR}/lib$(MPI64) -lmpi
     endif
 
 AOMP_GPU       ?= $(INSTALLED_GPU)
-AOMP_SANITIZER ?= $(AOMP_BUILD_SANITIZER)
 CC              = $(AOMP)/bin/clang
 OG11           ?= $(HOME)/git/og11/install
 OG11FLAGS      ?= -O3 -fopenmp -foffload=-march=$(AOMP_GPU) -I. -lgfortran -L $(OG11)/lib64
@@ -238,10 +237,10 @@ ifeq (gfx11,$(findstring gfx11,$(AOMP_GPU)))
  AOMP_WSIZE = -DWAVE_SIZE=32
 endif
 
-ifeq ($(AOMP_SANITIZER),)
-ASAN_FLAGS = -fsanitize=address -shared-libasan -g
-# For Sanitizer below environment variable should always be enabled by default
-AOMP_TARGET_FEATURES=:xnack+
+ifeq ($(AOMP_SANITIZER),1)
+  ASAN_FLAGS = -fsanitize=address -shared-libasan -g
+  #ASan requires xnack+ by default
+  AOMP_TARGET_FEATURES = :xnack+
 endif
 
 ifeq ($(AOMP_TARGET_FEATURES),)
diff --git a/test/smoke-asan/Makefile b/test/smoke-asan/Makefile
index 3577ffb8b..c97332d5a 100644
--- a/test/smoke-asan/Makefile
+++ b/test/smoke-asan/Makefile
@@ -3,9 +3,9 @@ include ../Makefile.defs
 TESTS_DIR = \
     hip-global-buffer-overflow \
     hip-heap-buffer-overflow \
-		hip-use-after-free \
-		omp-global-buffer-overflow \
-		omp-heap-buffer-overflow
+    hip-use-after-free \
+    omp-global-buffer-overflow \
+    omp-heap-buffer-overflow
 
 all:
 	@for test_dir in $(TESTS_DIR); do \
diff --git a/test/smoke/Makefile.rules b/test/smoke/Makefile.rules
index 821b8390c..3bf2e047f 100644
--- a/test/smoke/Makefile.rules
+++ b/test/smoke/Makefile.rules
@@ -125,6 +125,53 @@ else
 	@echo "  $(SKIP_RUN_SUPPORTED)"
 endif
 
+check-asan: $(TESTNAME)
+ifneq (,$(findstring $(GPU_W_FEATURES),$(SUPPORTED)))
+ifeq (,$(findstring $(GPU_W_FEATURES),$(UNSUPPORTED)))
+       path=`pwd`; \
+       base=`basename $$path`; \
+       ( \
+    flock -e 9 && echo  "" >> ../check-smoke-asan.txt; \
+    declare -A CmdStatus; \
+    pipeline="$(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1"; \
+    $(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1; \
+    pstat=($${PIPESTATUS[@]}); \
+    function GetPipedCmdStatus() { \
+      Cmd="$$2"; \
+      CmdRunIndex="$$3"; \
+      CmdRunIndexList=($$(echo $$1 | awk -F'|' '{for(i=1;i<=NF;i++) {gsub(/^[ \t]+|[ \t]+$$/, "", $$i); print $$i}}' | awk -v Cmd_Pattern="$$Cmd " '$$0~Cmd_Pattern {print NR}')); \
+      for CmdIndex in "$${CmdRunIndexList[@]}"; do \
+        if [ -v CmdStatus["$$Cmd"] ]; then \
+           ArrString=$${CmdStatus["$$Cmd"]}; \
+           ArrString+=",$${pstat[$$CmdIndex-1]}"; \
+           CmdStatus["$$Cmd"]=$${ArrString}; \
+        else \
+           NewArrString=$${pstat[$$CmdIndex-1]}; \
+           CmdStatus["$$Cmd"]=$${NewArrString}; \
+        fi; \
+      done; \
+      if [ -v CmdStatus[$$Cmd] ] && [ $$CmdRunIndex -gt 0 ];then \
+         IFS=',' read -ra CmdStatusList <<< "$${CmdStatus[$$Cmd]}"; \
+         return $${CmdStatusList["$$CmdRunIndex"-1]}; \
+      fi; \
+    }; \
+    GetPipedCmdStatus "$$pipeline" "$(TESTNAME)" "1"; \
+    test_status=$$?; \
+    GetPipedCmdStatus "$$pipeline" "FileCheck" "1"; \
+    filecheck_status=$$?; \
+         echo "$$test_status" > TEST_STATUS; \
+         echo "$$filecheck_status" > FILECHECK_STATUS; \
+         echo $$base $$test_num return code: $$test_status >> ../check-smoke-asan.txt; \
+         echo  "" >> ../check-smoke-asan.txt; \
+               if [ $$filecheck_status -eq 0 ]; then echo $$base $$test_num >> ../passing-tests.txt; \
+                                     else echo $$base $$test_num >> ../failing-tests.txt; fi; \
+       )9>../lockfile;
+else
+       @echo "  $(SKIP_RUN_UNSUPPORTED)"
+endif
+else
+       @echo "  $(SKIP_RUN_SUPPORTED)"
+endif
 
 # ----- Demo compile and link to object file
 ifneq ($(TESTNAME), $(findstring $(TESTNAME),$(TESTNAMES_ALL)))
@@ -249,7 +296,7 @@ run_sbin: sbin
 
 # Cleanup anything this makefile can create
 clean::
-	rm -f $(TESTNAME) $(TESTNAME).a llbin sbin obin *.i *.ii *.bc *.lk a.out-* *.ll *.s *.o *.log *.mod verify_output *.stb *.ilm *.cmod *.cmdx *.so $(TESTNAME)_og11 make-log.txt TEST_STATUS
+	rm -f $(TESTNAME) $(TESTNAME).a llbin sbin obin *.i *.ii *.bc *.lk a.out-* *.ll *.s *.o *.log *.mod verify_output *.stb *.ilm *.cmod *.cmdx *.so $(TESTNAME)_og11 make-log.txt TEST_STATUS FILECHECK_STATUS
 
 clean_log:
 	rm -f *.log
diff --git a/test/smoke/check_smoke.sh b/test/smoke/check_smoke.sh
index d52093c72..b49706a0a 100755
--- a/test/smoke/check_smoke.sh
+++ b/test/smoke/check_smoke.sh
@@ -242,12 +242,17 @@ if [ "$AOMP_PARALLEL_SMOKE" == 1 ]; then
         sem --jobs 4 --id def_sem -u 'base=$(basename $(pwd)); make check > /dev/null; if [ $? -ne 0 ]; then flock -e lockfile -c "echo $base: Make Failed >> ../make-fail.txt"; fi;'
       elif [ $base == "gpus" ]; then # Compile and link only test
         echo gpus is compile only!
+      elif [ "$AOMP_SANITIZER" == 1 ]; then
+        sem --jobs 4 --id def_sem -u 'make check-asan > /dev/null 2>&1'
       else
         sem --jobs 4 --id def_sem -u 'make check > /dev/null 2>&1'
       fi
       #---
       if [ -r "TEST_STATUS" ]; then
         test_status=`cat TEST_STATUS`
+        if [ "$AOMP_SANITIZER" == 1 ] && [ "$test_status" -ne 0 ]; then
+           test_status=`cat FILECHECK_STATUS`
+        fi
         if [ "$test_status" == "124" ]; then break; fi  # don't rerun timeouts
       fi
       run=$(($run+1))
@@ -332,6 +337,8 @@ for directory in $SMOKE_DIRS; do
     echo "$base" >> ../passing-tests.txt
   elif [ $base == 'printf_parallel_for_target' ] || [ $base == 'omp_places' ] || [ $base == 'pfspecifier' ] || [ $base == 'pfspecifier_str' ] ; then
     make verify-log
+  elif [ "$AOMP_SANITIZER" == 1 ]; then
+    make check-asan > /dev/null 2>&1
   else
     make check > /dev/null 2>&1
     # liba_bundled has an additional Makefile, that may fail on the make check
@@ -347,6 +354,9 @@ for directory in $SMOKE_DIRS; do
   #---
   if [ -r "TEST_STATUS" ]; then
     test_status=`cat TEST_STATUS`
+    if [ "$AOMP_SANITIZER" == 1 ] && [ "$test_status" -ne 0 ]; then
+       test_status=`cat FILECHECK_STATUS`
+    fi
     if [ "$test_status" == "124" ]; then break; fi      # don't rerun timeouts
   fi
   run=$(($run+1))

From c1da6ae10aaab951bd94d355beed3ae01f1e2cb9 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023@gmail.com>
Date: Tue, 16 Apr 2024 11:44:38 +0530
Subject: [PATCH 3/6] Minor Fixes

---
 test/Makefile.defs        |  1 +
 test/smoke/Makefile.rules | 10 +++++-----
 2 files changed, 6 insertions(+), 5 deletions(-)

diff --git a/test/Makefile.defs b/test/Makefile.defs
index 0b89e8f0d..4b8fb4c57 100644
--- a/test/Makefile.defs
+++ b/test/Makefile.defs
@@ -238,6 +238,7 @@ ifeq (gfx11,$(findstring gfx11,$(AOMP_GPU)))
 endif
 
 ifeq ($(AOMP_SANITIZER),1)
+  HSA_XNACK = 1
   ASAN_FLAGS = -fsanitize=address -shared-libasan -g
   #ASan requires xnack+ by default
   AOMP_TARGET_FEATURES = :xnack+
diff --git a/test/smoke/Makefile.rules b/test/smoke/Makefile.rules
index 3bf2e047f..b34094e15 100644
--- a/test/smoke/Makefile.rules
+++ b/test/smoke/Makefile.rules
@@ -159,11 +159,11 @@ ifeq (,$(findstring $(GPU_W_FEATURES),$(UNSUPPORTED)))
     test_status=$$?; \
     GetPipedCmdStatus "$$pipeline" "FileCheck" "1"; \
     filecheck_status=$$?; \
-         echo "$$test_status" > TEST_STATUS; \
-         echo "$$filecheck_status" > FILECHECK_STATUS; \
-         echo $$base $$test_num return code: $$test_status >> ../check-smoke-asan.txt; \
-         echo  "" >> ../check-smoke-asan.txt; \
-               if [ $$filecheck_status -eq 0 ]; then echo $$base $$test_num >> ../passing-tests.txt; \
+    echo "$$test_status" > TEST_STATUS; \
+    echo "$$filecheck_status" > FILECHECK_STATUS; \
+    echo $$base $$test_num return code: $$test_status >> ../check-smoke-asan.txt; \
+    echo  "" >> ../check-smoke-asan.txt; \
+    if [ $$filecheck_status -eq 0 ]; then echo $$base $$test_num >> ../passing-tests.txt; \
                                      else echo $$base $$test_num >> ../failing-tests.txt; fi; \
        )9>../lockfile;
 else

From 3ed4b09c7dc524afcffd63ffc821739de1b605a7 Mon Sep 17 00:00:00 2001
From: Amit Pandey <amit.pandey@amd.com>
Date: Wed, 17 Apr 2024 15:53:28 +0530
Subject: [PATCH 4/6] Addressing Comments.

1. Add comments related to documentation for check-asan target.
2. Ensure that check-asan target should only be executed for smoke-asan
   tests.
---
 test/smoke-asan/Makefile  | 14 +++---
 test/smoke/Makefile.rules | 98 +++++++++++++++++++++++----------------
 test/smoke/check_smoke.sh |  5 +-
 3 files changed, 68 insertions(+), 49 deletions(-)

diff --git a/test/smoke-asan/Makefile b/test/smoke-asan/Makefile
index c97332d5a..5b2a2c3fe 100644
--- a/test/smoke-asan/Makefile
+++ b/test/smoke-asan/Makefile
@@ -23,13 +23,13 @@ run run_obin run_sbin run_llbin clean clean_log llbin sbin obin:
 	  $(MAKE) -C $$test_dir $@; \
 	done
 
-check:
-	 @for test_dir in $(TESTS_DIR); do \
-          echo $$nnn; \
-          test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
-          echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir $@"; \
-          $(MAKE) -C $$test_dir $@; \
-        done
+check-asan:
+	@for test_dir in $(TESTS_DIR); do \
+	  echo $$nnn; \
+	  test_name=`grep "TESTNAME *=" $$test_dir/Makefile | sed "s/.*= *//"`; \
+	  echo "TEST_DIR: $$test_dir\tTEST_NAME: $$test_name\tMAKE: $(MAKE) -C $$test_dir $@"; \
+	  $(MAKE) -C $$test_dir $@; \
+	done
 
 .ll .ll.s .ll.o .s .s.o .o:
 	@for test_dir in $(TESTS_DIR); do \
diff --git a/test/smoke/Makefile.rules b/test/smoke/Makefile.rules
index b34094e15..77e6342da 100644
--- a/test/smoke/Makefile.rules
+++ b/test/smoke/Makefile.rules
@@ -128,51 +128,69 @@ endif
 check-asan: $(TESTNAME)
 ifneq (,$(findstring $(GPU_W_FEATURES),$(SUPPORTED)))
 ifeq (,$(findstring $(GPU_W_FEATURES),$(UNSUPPORTED)))
-       path=`pwd`; \
-       base=`basename $$path`; \
-       ( \
-    flock -e 9 && echo  "" >> ../check-smoke-asan.txt; \
-    declare -A CmdStatus; \
-    pipeline="$(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1"; \
-    $(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1; \
-    pstat=($${PIPESTATUS[@]}); \
-    function GetPipedCmdStatus() { \
-      Cmd="$$2"; \
-      CmdRunIndex="$$3"; \
-      CmdRunIndexList=($$(echo $$1 | awk -F'|' '{for(i=1;i<=NF;i++) {gsub(/^[ \t]+|[ \t]+$$/, "", $$i); print $$i}}' | awk -v Cmd_Pattern="$$Cmd " '$$0~Cmd_Pattern {print NR}')); \
-      for CmdIndex in "$${CmdRunIndexList[@]}"; do \
-        if [ -v CmdStatus["$$Cmd"] ]; then \
-           ArrString=$${CmdStatus["$$Cmd"]}; \
-           ArrString+=",$${pstat[$$CmdIndex-1]}"; \
-           CmdStatus["$$Cmd"]=$${ArrString}; \
-        else \
-           NewArrString=$${pstat[$$CmdIndex-1]}; \
-           CmdStatus["$$Cmd"]=$${NewArrString}; \
-        fi; \
-      done; \
-      if [ -v CmdStatus[$$Cmd] ] && [ $$CmdRunIndex -gt 0 ];then \
-         IFS=',' read -ra CmdStatusList <<< "$${CmdStatus[$$Cmd]}"; \
-         return $${CmdStatusList["$$CmdRunIndex"-1]}; \
-      fi; \
-    }; \
-    GetPipedCmdStatus "$$pipeline" "$(TESTNAME)" "1"; \
-    test_status=$$?; \
-    GetPipedCmdStatus "$$pipeline" "FileCheck" "1"; \
-    filecheck_status=$$?; \
-    echo "$$test_status" > TEST_STATUS; \
-    echo "$$filecheck_status" > FILECHECK_STATUS; \
-    echo $$base $$test_num return code: $$test_status >> ../check-smoke-asan.txt; \
-    echo  "" >> ../check-smoke-asan.txt; \
-    if [ $$filecheck_status -eq 0 ]; then echo $$base $$test_num >> ../passing-tests.txt; \
-                                     else echo $$base $$test_num >> ../failing-tests.txt; fi; \
-       )9>../lockfile;
+	path=`pwd`; \
+	base=`basename $$path`; \
+	( \
+	  flock -e 9 && echo  "" >> ../check-smoke-asan.txt; \
+	  # CmdStatus: Associative Container which keeps track of execution status of any command in the pipeline. \
+	  declare -A CmdStatus; \
+	  # GetPipedCmdStatus: Function which on demand populates the CmdStatus container with execution status values of each command executing in a given input pipeline. \
+	  # param_1: pipeline as string. \
+	  # param_2: pipestatus computed from PIPESTATUS as array. \
+	  # param_3: Cmd Name as string. \
+	  # param_4: CmdRunIndex of Cmd as positive integer. \
+	  function GetPipedCmdStatus() { \
+	  # pstat: Piped command status list. \
+	    local -n pstat=$$2; \
+	  # Cmd: Command Name \
+	    Cmd="$$3"; \
+	  # CmdRunIndex: Index value of Cmd executing in the pipeline. \
+	  # Ex. pipeline = "Cmd1 | Cmd2 | Cmd3 | Cmd1 | Cmd4 | Cmd3" \
+	  # CmdRunIndex of Cmd1(First Time): 1 \
+	  # CmdRunIndex of Cmd2(First Time): 1 \
+	  # CmdRunIndex of Cmd1(Second Time): 2 \
+	    CmdRunIndex="$$4"; \
+	  # CmdRunIndexList: Index List of each Cmd. \
+	  # Ex. pipeline = "Cmd1 | Cmd2 | Cmd1 | Cmd3" \
+	  #     CmdRunIndexList of Cmd1: (1,3) \
+	    CmdRunIndexList=($$(echo $$1 | awk -F'|' '{for(i=1;i<=NF;i++) {gsub(/^[ \t]+|[ \t]+$$/, "", $$i); print $$i}}' | awk -v Cmd_Pattern="$$Cmd " '$$0~Cmd_Pattern {print NR}')); \
+	    for CmdIndex in "$${CmdRunIndexList[@]}"; do \
+	      if [ -v CmdStatus["$$Cmd"] ]; then \
+	         ArrString=$${CmdStatus["$$Cmd"]}; \
+	         ArrString+=",$${pstat[$$CmdIndex-1]}"; \
+	         CmdStatus["$$Cmd"]=$${ArrString}; \
+	      else \
+	         NewArrString=$${pstat[$$CmdIndex-1]}; \
+	         CmdStatus["$$Cmd"]=$${NewArrString}; \
+	      fi; \
+	    done; \
+	    if [ -v CmdStatus[$$Cmd] ] && [ $$CmdRunIndex -gt 0 ];then \
+	       IFS=',' read -ra CmdStatusList <<< "$${CmdStatus[$$Cmd]}"; \
+	       return $${CmdStatusList["$$CmdRunIndex"-1]}; \
+	    fi; \
+	  }; \
+	  pipeline="$(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1"; \
+	  $(RUNENV) $(SMOKE_TIMEOUT) $(RUNPROF) $(RUNPROF_FLAGS) $(CHECK_COMMAND) > /dev/null 2>&1; \
+	  pipestatus=($${PIPESTATUS[@]}); \
+	  GetPipedCmdStatus "$$pipeline" "pipestatus" "$(TESTNAME)" "1"; \
+	  test_status=$$?; \
+	  GetPipedCmdStatus "$$pipeline" "pipestatus" "FileCheck" "1"; \
+	  filecheck_status=$$?; \
+	  echo "$$test_status" > TEST_STATUS; \
+	  echo "$$filecheck_status" > FILECHECK_STATUS; \
+	  echo $$base $$test_num return code: $$test_status >> ../check-smoke-asan.txt; \
+	  echo  "" >> ../check-smoke-asan.txt; \
+	  if [ $$filecheck_status -eq 0 ]; then echo $$base $$test_num >> ../passing-tests.txt; \
+	                                   else echo $$base $$test_num >> ../failing-tests.txt; fi; \
+	)9>../lockfile;
 else
-       @echo "  $(SKIP_RUN_UNSUPPORTED)"
+	@echo "  $(SKIP_RUN_UNSUPPORTED)"
 endif
 else
-       @echo "  $(SKIP_RUN_SUPPORTED)"
+	@echo "  $(SKIP_RUN_SUPPORTED)"
 endif
 
+
 # ----- Demo compile and link to object file
 ifneq ($(TESTNAME), $(findstring $(TESTNAME),$(TESTNAMES_ALL)))
 .PHONY: $(TESTNAME).o
diff --git a/test/smoke/check_smoke.sh b/test/smoke/check_smoke.sh
index b49706a0a..a1d40760b 100755
--- a/test/smoke/check_smoke.sh
+++ b/test/smoke/check_smoke.sh
@@ -90,6 +90,7 @@ cleanup(){
 script_dir=$(dirname "$0")
 pushd $script_dir
 path=$(pwd)
+script_dir_name=$(basename "$path")
 
 SMOKE_DIRS=${SMOKE_DIRS:-./*/}  # test directories to run
 SMOKE_LRUN=${SMOKE_LRUN:-1}     # number of times to run test list
@@ -242,7 +243,7 @@ if [ "$AOMP_PARALLEL_SMOKE" == 1 ]; then
         sem --jobs 4 --id def_sem -u 'base=$(basename $(pwd)); make check > /dev/null; if [ $? -ne 0 ]; then flock -e lockfile -c "echo $base: Make Failed >> ../make-fail.txt"; fi;'
       elif [ $base == "gpus" ]; then # Compile and link only test
         echo gpus is compile only!
-      elif [ "$AOMP_SANITIZER" == 1 ]; then
+      elif [ "$AOMP_SANITIZER" == 1 ] && [ "$script_dir_name" == "smoke-asan" ]; then
         sem --jobs 4 --id def_sem -u 'make check-asan > /dev/null 2>&1'
       else
         sem --jobs 4 --id def_sem -u 'make check > /dev/null 2>&1'
@@ -337,7 +338,7 @@ for directory in $SMOKE_DIRS; do
     echo "$base" >> ../passing-tests.txt
   elif [ $base == 'printf_parallel_for_target' ] || [ $base == 'omp_places' ] || [ $base == 'pfspecifier' ] || [ $base == 'pfspecifier_str' ] ; then
     make verify-log
-  elif [ "$AOMP_SANITIZER" == 1 ]; then
+  elif [ "$AOMP_SANITIZER" == 1 ] && [ "$script_dir_name" == "smoke-asan" ]; then
     make check-asan > /dev/null 2>&1
   else
     make check > /dev/null 2>&1

From 635627721f9751ec097b8f48e4b8905326e136cf Mon Sep 17 00:00:00 2001
From: Amit Pandey <amit.pandey@amd.com>
Date: Mon, 22 Apr 2024 16:22:48 +0530
Subject: [PATCH 5/6] Add omp-heap-use-after-free test.

---
 test/Makefile.defs                            |  3 ++-
 test/smoke-asan/Makefile                      |  3 ++-
 .../hip-global-buffer-overflow/Makefile       |  2 +-
 .../hip-heap-buffer-overflow/Makefile         |  2 +-
 test/smoke-asan/hip-use-after-free/Makefile   |  2 +-
 .../omp-global-buffer-overflow/Makefile       |  2 +-
 .../omp-global-buffer-overflow.cpp            |  2 +-
 .../omp-heap-buffer-overflow/Makefile         |  2 +-
 .../omp-heap-buffer-overflow.cpp              |  2 +-
 .../omp-heap-use-after-free/Makefile          | 21 ++++++++++++++++++
 .../omp-heap-use-after-free.cpp               | 22 +++++++++++++++++++
 11 files changed, 54 insertions(+), 9 deletions(-)
 create mode 100644 test/smoke-asan/omp-heap-use-after-free/Makefile
 create mode 100644 test/smoke-asan/omp-heap-use-after-free/omp-heap-use-after-free.cpp

diff --git a/test/Makefile.defs b/test/Makefile.defs
index 4b8fb4c57..c9b45898a 100644
--- a/test/Makefile.defs
+++ b/test/Makefile.defs
@@ -26,7 +26,8 @@ else
   # target IDs that support OpenMP 'requires unified_shared_memory'
   # as defined in https://llvm.org/docs/AMDGPUUsage.html#amdgpu-processor-table
   # with xnack target feature supported.
-  SUPPORTS_USM= gfx90a, gfx90a:xnack+,\
+  SUPPORTS_USM= gfx908, gfx908:xnack+,\
+                gfx90a, gfx90a:xnack+,\
                 gfx90c, gfx90c:xnack+,\
                 gfx940, gfx940:xnack+,\
                 gfx941, gfx941:xnack+,\
diff --git a/test/smoke-asan/Makefile b/test/smoke-asan/Makefile
index 5b2a2c3fe..b916854c4 100644
--- a/test/smoke-asan/Makefile
+++ b/test/smoke-asan/Makefile
@@ -5,7 +5,8 @@ TESTS_DIR = \
     hip-heap-buffer-overflow \
     hip-use-after-free \
     omp-global-buffer-overflow \
-    omp-heap-buffer-overflow
+    omp-heap-buffer-overflow \
+    omp-heap-use-after-free
 
 all:
 	@for test_dir in $(TESTS_DIR); do \
diff --git a/test/smoke-asan/hip-global-buffer-overflow/Makefile b/test/smoke-asan/hip-global-buffer-overflow/Makefile
index 2c94934fc..64a8df7e2 100644
--- a/test/smoke-asan/hip-global-buffer-overflow/Makefile
+++ b/test/smoke-asan/hip-global-buffer-overflow/Makefile
@@ -22,7 +22,7 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
 
 #-ccc-print-phases
 #"-\#\#\#"
diff --git a/test/smoke-asan/hip-heap-buffer-overflow/Makefile b/test/smoke-asan/hip-heap-buffer-overflow/Makefile
index 6c1a9c4fd..b5679cac3 100644
--- a/test/smoke-asan/hip-heap-buffer-overflow/Makefile
+++ b/test/smoke-asan/hip-heap-buffer-overflow/Makefile
@@ -22,7 +22,7 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
 
 #-ccc-print-phases
 #"-\#\#\#"
diff --git a/test/smoke-asan/hip-use-after-free/Makefile b/test/smoke-asan/hip-use-after-free/Makefile
index afd6bc684..8b8013bcf 100644
--- a/test/smoke-asan/hip-use-after-free/Makefile
+++ b/test/smoke-asan/hip-use-after-free/Makefile
@@ -22,7 +22,7 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
 
 #-ccc-print-phases
 #"-\#\#\#"
diff --git a/test/smoke-asan/omp-global-buffer-overflow/Makefile b/test/smoke-asan/omp-global-buffer-overflow/Makefile
index 6521eba03..177d0184d 100644
--- a/test/smoke-asan/omp-global-buffer-overflow/Makefile
+++ b/test/smoke-asan/omp-global-buffer-overflow/Makefile
@@ -14,7 +14,7 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
 #-ccc-print-phases
 #"-\#\#\#"
 
diff --git a/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp b/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
index 788f33701..d4f0fc2ca 100644
--- a/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
+++ b/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp
@@ -19,4 +19,4 @@ int main(int argc, char *argv[]) {
 /// CHECK:=================================================================
 /// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: global-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
 /// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
-/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_fd00_68a0014_main_l11 at {{.*}}aomp/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp:13:13
+/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_{{.*}} at {{.*}}aomp/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp:13:13
diff --git a/test/smoke-asan/omp-heap-buffer-overflow/Makefile b/test/smoke-asan/omp-heap-buffer-overflow/Makefile
index 399089885..221ab4e0d 100644
--- a/test/smoke-asan/omp-heap-buffer-overflow/Makefile
+++ b/test/smoke-asan/omp-heap-buffer-overflow/Makefile
@@ -14,7 +14,7 @@ CLANG        ?= clang++
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
 
-SUPPORTED    = gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
 #-ccc-print-phases
 #"-\#\#\#"
 
diff --git a/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp b/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp
index eb54d8a84..dfdc1ad13 100644
--- a/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp
+++ b/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp
@@ -15,4 +15,4 @@ int main(int argc, char *argv[]) {
 /// CHECK:=================================================================
 /// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: heap-buffer-overflow on amdgpu device 0 at pc [[PC:.*]]
 /// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
-/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_fd00_68a0017_main_l7 at {{.*}}aomp/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp:9:11
+/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_{{.*}} at {{.*}}aomp/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp:9:11
diff --git a/test/smoke-asan/omp-heap-use-after-free/Makefile b/test/smoke-asan/omp-heap-use-after-free/Makefile
new file mode 100644
index 000000000..26c62caf3
--- /dev/null
+++ b/test/smoke-asan/omp-heap-use-after-free/Makefile
@@ -0,0 +1,21 @@
+include ../../Makefile.defs
+
+TESTNAME     = omp-heap-use-after-free
+TESTSRC_MAIN = omp-heap-use-after-free.cpp
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+HSA_XNACK ?= 1
+RUNENV      += HSA_XNACK=$(HSA_XNACK)
+
+RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) --strict-whitespace --match-full-lines $(TESTSRC_MAIN)
+
+CLANG        ?= clang++
+OMP_BIN      = $(AOMP)/bin/$(CLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+
+SUPPORTED    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
diff --git a/test/smoke-asan/omp-heap-use-after-free/omp-heap-use-after-free.cpp b/test/smoke-asan/omp-heap-use-after-free/omp-heap-use-after-free.cpp
new file mode 100644
index 000000000..2fce0ea1a
--- /dev/null
+++ b/test/smoke-asan/omp-heap-use-after-free/omp-heap-use-after-free.cpp
@@ -0,0 +1,22 @@
+#include <omp.h>
+
+int main(int argc, char *argv[]) {
+  const unsigned long int N = 10000;
+
+  float *buffer = (float *)omp_target_alloc(N * sizeof(float), 0);
+
+  omp_target_free(buffer, 0);
+#pragma omp target teams num_teams(2) is_device_ptr(buffer)
+  {
+#pragma omp parallel for
+    for (unsigned long int i = 0; i < N; ++i) {
+      buffer[i + 1] = i;
+    }
+  }
+  return 0;
+}
+
+/// CHECK:=================================================================
+/// CHECK-NEXT:=={{[0-9]+}}==ERROR: AddressSanitizer: heap-use-after-free on amdgpu device 0 at pc [[PC:.*]]
+/// CHECK-NEXT:WRITE of size 4 in workgroup id ({{[0-9]+}},0,0)
+/// CHECK-NEXT:  #0 [[PC]] in __omp_offloading_{{.*}} at {{.*}}aomp/test/smoke-asan/omp-heap-use-after-free/omp-heap-use-after-free.cpp:13:21

From af8447293062236553a8ea59f70e37aa360369c2 Mon Sep 17 00:00:00 2001
From: Amit Pandey <amit.pandey@amd.com>
Date: Thu, 25 Apr 2024 15:48:21 +0530
Subject: [PATCH 6/6] Add AOMP ASan Smoke testing using command line arguments.

Usage: ./run_rocm_test.sh -a
---
 bin/run_rocm_test.sh      | 25 ++++++++++++++++++++++---
 test/smoke/Makefile.rules |  2 +-
 2 files changed, 23 insertions(+), 4 deletions(-)

diff --git a/bin/run_rocm_test.sh b/bin/run_rocm_test.sh
index 0c9dcb124..738c53346 100755
--- a/bin/run_rocm_test.sh
+++ b/bin/run_rocm_test.sh
@@ -14,6 +14,12 @@
 # we need to see 1 device only, babelstream in particular.
 export ROCR_VISIBLE_DEVICES=0
 
+# Enable AMDGPU Sanitizer Testing
+if [ "$1" == "-a" ]; then
+  export AOMP_SANITIZER=1
+  export LD_LIBRARY_PATH=$ROCM_INSTALL_PATH/llvm/lib/asan:$ROCM_INSTALL_PATH/lib/asan:$LD_LIBRARY_PATH
+fi
+
 if [ -e /usr/sbin/lspci ]; then
   lspci_loc=/usr/sbin/lspci
 else
@@ -37,13 +43,13 @@ if [ $ISVIRT -eq 1 ] ; then
 SKIP_USM=1
 export SKIP_USM=1
 export HSA_XNACK=${HSA_XNACK:-0}
-SUITE_LIST=${SUITE_LIST:-"examples smoke-limbo smoke omp5 openmpapps ovo sollve babelstream fortran-babelstream"}
+SUITE_LIST=${SUITE_LIST:-"examples smoke-limbo smoke smoke-asan omp5 openmpapps ovo sollve babelstream fortran-babelstream"}
 blockinglist="examples_fortran examples_openmp smoke openmpapps sollve45 sollve50 babelstream"
 else
-SUITE_LIST=${SUITE_LIST:-"examples smoke-limbo smoke omp5 openmpapps LLNL nekbone ovo sollve babelstream fortran-babelstream"}
+SUITE_LIST=${SUITE_LIST:-"examples smoke-limbo smoke smoke-asan omp5 openmpapps LLNL nekbone ovo sollve babelstream fortran-babelstream"}
 blockinglist="examples_fortran examples_openmp smoke openmpapps sollve45 sollve50 babelstream"
 fi
-EPSDB_LIST=${EPSDB_LIST:-"examples smoke-limbo smoke omp5 openmpapps LLNL nekbone ovo sollve babelstream fortran-babelstream"}
+EPSDB_LIST=${EPSDB_LIST:-"examples smoke-limbo smoke smoke-asan omp5 openmpapps LLNL nekbone ovo sollve babelstream fortran-babelstream"}
 
 export AOMP_USE_CCACHE=0
 
@@ -546,6 +552,19 @@ function smoke(){
   copyresults smoke "$aompdir"/test/smoke
 }
 
+function smoke-asan(){
+  # Smoke-ASan
+  if [ "$AOMP_SANITIZER" == 1 ]; then
+    mkdir -p "$resultsdir"/smoke-asan
+    cd "$aompdir"/test/smoke-asan
+    HIP_PATH="" AOMP_PARALLEL_SMOKE=1 CLEANUP=0 AOMPHIP=$AOMPROCM ./check_smoke_asan.sh
+    checkrc $?
+    copyresults smoke-asan "$aompdir"/test/smoke-asan
+  else
+    echo "Skipping smoke-asan."
+  fi
+}
+
 SMOKE_FAILS=${SMOKE_FAILS:-1}
 function smokefails(){
   # Smoke-fails
diff --git a/test/smoke/Makefile.rules b/test/smoke/Makefile.rules
index 77e6342da..c00ca2955 100644
--- a/test/smoke/Makefile.rules
+++ b/test/smoke/Makefile.rules
@@ -140,7 +140,7 @@ ifeq (,$(findstring $(GPU_W_FEATURES),$(UNSUPPORTED)))
 	  # param_3: Cmd Name as string. \
 	  # param_4: CmdRunIndex of Cmd as positive integer. \
 	  function GetPipedCmdStatus() { \
-	  # pstat: Piped command status list. \
+	  # pstat: Pipeline status list. \
 	    local -n pstat=$$2; \
 	  # Cmd: Command Name \
 	    Cmd="$$3"; \