Add Smoke tests for ASan.

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

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+,\
@@ -237,6 +238,13 @@ ifeq (gfx11,$(findstring gfx11,$(AOMP_GPU)))
  AOMP_WSIZE = -DWAVE_SIZE=32
 endif
 
+ifeq ($(AOMP_SANITIZER),1)
+  HSA_XNACK = 1
+  ASAN_FLAGS = -fsanitize=address -shared-libasan -g
+  #ASan requires xnack+ by default
+  AOMP_TARGET_FEATURES = :xnack+
+endif
+
 ifeq ($(AOMP_TARGET_FEATURES),)
 GPU_W_FEATURES = $(AOMP_GPU)
 else
@@ -247,9 +255,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

test/smoke-asan/Makefile

@@ -0,0 +1,90 @@
+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 \
+    omp-heap-use-after-free
+
+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-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 \
+	  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

test/smoke-asan/Makefile.rules

@@ -0,0 +1 @@
+../smoke/Makefile.rules

test/smoke-asan/check_smoke_asan.sh

@@ -0,0 +1 @@
+../smoke/check_smoke.sh

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    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules

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

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    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules

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

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    = gfx908:xnack+,gfx90a:xnack+,gfx940:xnack+,gfx941:xnack+,gfx942:xnack+
+
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules

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