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/Makefile.defs b/test/Makefile.defs index e12278e9d..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+,\ @@ -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 diff --git a/test/smoke-asan/Makefile b/test/smoke-asan/Makefile new file mode 100644 index 000000000..b916854c4 --- /dev/null +++ b/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 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..64a8df7e2 --- /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 = gfx908:xnack+,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 + +#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..b5679cac3 --- /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 = gfx908:xnack+,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 + +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..8b8013bcf --- /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 = gfx908:xnack+,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 + +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..177d0184d --- /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 = gfx908:xnack+,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..d4f0fc2ca --- /dev/null +++ b/test/smoke-asan/omp-global-buffer-overflow/omp-global-buffer-overflow.cpp @@ -0,0 +1,22 @@ +#include +#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_{{.*}} 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..221ab4e0d --- /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 = gfx908:xnack+,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..dfdc1ad13 --- /dev/null +++ b/test/smoke-asan/omp-heap-buffer-overflow/omp-heap-buffer-overflow.cpp @@ -0,0 +1,18 @@ +#include + +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_{{.*}} 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 + +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 diff --git a/test/smoke/Makefile.rules b/test/smoke/Makefile.rules index 821b8390c..c00ca2955 100644 --- a/test/smoke/Makefile.rules +++ b/test/smoke/Makefile.rules @@ -125,6 +125,71 @@ 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; \ + # 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: Pipeline 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)" +endif +else + @echo " $(SKIP_RUN_SUPPORTED)" +endif + # ----- Demo compile and link to object file ifneq ($(TESTNAME), $(findstring $(TESTNAME),$(TESTNAMES_ALL))) @@ -249,7 +314,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..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,12 +243,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 ] && [ "$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' 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 +338,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 ] && [ "$script_dir_name" == "smoke-asan" ]; 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 +355,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))