Merge branch 'aomp-epsdb' into aomp-epsdb-mainline

Change-Id: Ia590b2318f03fe2b366e80948e4f20d160728fc5
ROCm · Aug 10, 2022 · 4590916 · 4590916
2 parents 6f1bb20 + ec492be
commit 4590916
Show file tree

Hide file tree

Showing 8 changed files with 564 additions and 3 deletions.
diff --git a/bin/build_openmp.sh b/bin/build_openmp.sh
@@ -82,6 +82,7 @@ GFXSEMICOLONS=`echo $GFXLIST | tr ' ' ';' `
 #COMMON_CMAKE_OPTS="#-DOPENMP_TEST_C_COMPILER=$AOMP/bin/clang
 #-DOPENMP_TEST_CXX_COMPILER=$AOMP/bin/clang++
 
+# FIXME: Remove CMAKE_CXX_FLAGS and CMAKE_C_FLAGS when AFAR uses 5.3 ROCr.
 COMMON_CMAKE_OPTS="-DDEVICELIBS_ROOT=$DEVICELIBS_ROOT
 -DOPENMP_ENABLE_LIBOMPTARGET=1
 -DOPENMP_ENABLE_LIBOMPTARGET_HSA=1
@@ -97,7 +98,9 @@ COMMON_CMAKE_OPTS="-DDEVICELIBS_ROOT=$DEVICELIBS_ROOT
 -DCMAKE_C_COMPILER=$OUT_DIR/llvm/bin/clang
 -DCMAKE_CXX_COMPILER=$OUT_DIR/llvm/bin/clang++
 -DOPENMP_TEST_C_COMPILER=$OUT_DIR/llvm/bin/clang
--DOPENMP_TEST_CXX_COMPILER=$OUT_DIR/llvm/bin/clang++"
+-DOPENMP_TEST_CXX_COMPILER=$OUT_DIR/llvm/bin/clang++
+-DCMAKE_CXX_FLAGS=-I$ROCM_DIR/include
+-DCMAKE_C_FLAGS=-I$ROCM_DIR/include"
 
 if [ "$AOMP_STANDALONE_BUILD" == 0 ]; then
   COMMON_CMAKE_OPTS="$COMMON_CMAKE_OPTS -DENABLE_DEVEL_PACKAGE=ON -DENABLE_RUN_PACKAGE=ON"
@@ -183,8 +186,9 @@ if [ "$1" != "nocmake" ] && [ "$1" != "install" ] ; then
          echo ${AOMP_CMAKE} $MYCMAKEOPTS  $BUILD_DIR/$AOMP_PROJECT_REPO_NAME/openmp
          env "$@" ${AOMP_CMAKE} $MYCMAKEOPTS  $BUILD_DIR/$AOMP_PROJECT_REPO_NAME/openmp
       else
-         echo ${AOMP_CMAKE} $MYCMAKEOPTS  $AOMP_REPOS/../$AOMP_PROJECT_REPO_NAME/openmp
-         env "$@" ${AOMP_CMAKE} $MYCMAKEOPTS  $AOMP_REPOS/../$AOMP_PROJECT_REPO_NAME/openmp
+         # FIXME: Remove CMAKE_CXX_FLAGS and CMAKE_C_FLAGS when AFAR uses 5.3 ROCr.
+         echo ${AOMP_CMAKE} $MYCMAKEOPTS -DCMAKE_CXX_FLAGS="-g -I$ROCM_DIR/include" -DCMAKE_C_FLAGS="-g -I$ROCM_DIR/include" $AOMP_REPOS/../$AOMP_PROJECT_REPO_NAME/openmp
+         env "$@" ${AOMP_CMAKE} $MYCMAKEOPTS -DCMAKE_CXX_FLAGS="-g -I$ROCM_DIR/include" -DCMAKE_C_FLAGS="-g -I$ROCM_DIR/include" $AOMP_REPOS/../$AOMP_PROJECT_REPO_NAME/openmp
       fi
       if [ $? != 0 ] ; then 
          echo "ERROR openmp debug cmake failed. Cmake flags"

diff --git a/examples/tools/asan/README.md b/examples/tools/asan/README.md
@@ -0,0 +1,31 @@
+Address Sanitizer(ASan) is a memory error detector tool utilized by applications to detect various errors ranging from spatial issues like out-of-bound access to temporal issues like use-after-free.
+
+* **Features Supported On Host Platfrom(Target x86_64).**
+   1. Use after free
+   2. Buffer Overflows
+        - Heap buffer overflow
+        - Stack buffer overflow
+        - Global buffer overflow
+   3. Use after return
+   4. Use after scope
+   5. Initialization order bugs
+
+ * **Features Supported on AMDGPU Platform(amdgcn-amd-amdhsa)**
+   1. Heap buffer overflow
+   2. Global buffer overflow
+
+Requirements
+========================================================================================================================================
+
+* **Software(Kernel/OS) Requirements**
+
+    - **Unified Memory Support(HMM Kernel)**
+        1. This feature requires Linux Kernel versions greater than 5.14.
+        2. This feature also requires latest KFD driver packaged in ROCm stack.
+        3. Unified memory support can be tested with applications compiled with xnack capability.
+
+    - **XNACK Capability**
+        1. Xnack replay enabled mode compiled binaries on execution indicates that runtime can handle page faults gracefully.So,if any page faults occur on gpu then at runtime a retry of that memory access happens.
+              > xnack+ --offload-arch=gfx908:xnack+
+        2. Xnack replay disabled mode compiled binaries on indicates that runtime can't handle page faults on GPU.So, developer should write offloading kernels with caution that it should not create any page faults on GPU.
+              > xnack- with –offload-arch=gfx908:xnack-
diff --git a/examples/tools/asan/global_buffer_overflow/README.md b/examples/tools/asan/global_buffer_overflow/README.md
@@ -0,0 +1,71 @@
+**Global Buffer Overflow** Demonstration on AMDGPU with applications written in HIP/OpenMP.
+========================================================================================================================================
+
+> ## **HIP** ##
+>
+> ```CPP
+>
+> __global__ void vecAdd(int *A,int *B,int *C,int N){
+> int i = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+> if(i < N){
+>   C[i+10] = A[i] + B[i];
+> }
+> }
+>
+>```
+
+>
+> ## **OpenMP** ##
+>
+>
+>```CPP
+>
+>#pragma omp target parallel for
+>for(int i = 0; i < N; i++){
+>    C[i+10] = A[i] + B[i];
+>}
+>
+>=================================================================
+>==3472068==ERROR: AddressSanitizer: global-buffer-overflow on amdgpu device 0 at pc 0x7f150ae13cb4
+>WRITE of size 4 in workgroup id (0,0,0)
+>  #0 0x7f150ae13cb4 in __omp_outlined__ at /home/ampandey/device-asan/openmp/vecadd-GBO.cpp:32:8
+>
+>Thread ids and accessed addresses:
+>90 : 0x7f150ae7a850 91 : 0x7f150ae7a854 92 : 0x7f150ae7a858 93 : 0x7f150ae7a85c 94 : 0x7f150ae7a860 95 : 0x7f150ae7a864 96 : 0x7f150ae7a868 97 : 0x7f150ae7a86c
+>98 : 0x7f150ae7a870
+>
+>Address 0x7f150ae7a850 is a wild pointer inside of access range of size 0x000000000004.
+>Shadow bytes around the buggy address:
+>   0x0fe3215c74b0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>   0x0fe3215c74c0: 00 00 00 00 00 00 00 00 00 00 f9 f9 f9 f9 f9 f9
+>   0x0fe3215c74d0: f9 f9 f9 f9 f9 f9 f9 f9 00 00 00 00 00 00 00 00
+>   0x0fe3215c74e0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>   0x0fe3215c74f0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+> =>0x0fe3215c7500: 00 00 00 00 00 00 00 00 00 00[f9]f9 f9 f9 f9 f9
+>   0x0fe3215c7510: f9 f9 f9 f9 f9 f9 f9 f9 00 00 00 00 00 00 00 00
+>   0x0fe3215c7520: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>   0x0fe3215c7530: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>   0x0fe3215c7540: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>   0x0fe3215c7550: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
+>Shadow byte legend (one shadow byte represents 8 application bytes):
+>  Addressable:           00
+>  Partially addressable: 01 02 03 04 05 06 07
+>  Heap left redzone:       fa
+>  Freed heap region:       fd
+>  Stack left redzone:      f1
+>  Stack mid redzone:       f2
+>  Stack right redzone:     f3
+>  Stack after return:      f5
+>  Stack use after scope:   f8
+>  Global redzone:          f9
+>  Global init order:       f6
+>  Poisoned by user:        f7
+>  Container overflow:      fc
+>  Array cookie:            ac
+>  Intra object redzone:    bb
+>  ASan internal:           fe
+>  Left alloca redzone:     ca
+>  Right alloca redzone:    cb
+>==3472068==ABORTING
+'''
+>
diff --git a/examples/tools/asan/global_buffer_overflow/openmp/Makefile b/examples/tools/asan/global_buffer_overflow/openmp/Makefile
@@ -0,0 +1,123 @@
+#-----------------------------------------------------------------------
+#
+#  Makefile: Compile OpenMP test case with ASan Flags.
+#
+#  Run "make help" to see how to use this Makefile
+#
+#-----------------------------------------------------------------------
+# MIT License
+# Copyright (c) 2017 Advanced Micro Devices, Inc. All Rights Reserved.
+#
+# Permission is hereby granted, free of charge, to any person
+# obtaining a copy of this software and associated documentation
+# files (the "Software"), to deal in the Software without
+# restriction, including without limitation the rights to use, copy,
+# modify, merge, publish, distribute, sublicense, and/or sell copies
+# of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be
+# included in all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+# EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+# MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+# NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS
+# BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN
+# ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
+# CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+TESTNAME =vecadd-GBO
+FILETYPE =cpp
+
+ROCM_LLVM ?= /opt/rocm/llvm
+
+ifeq ($(LLVM_GPU),)
+INSTALLED_GPU  = $(shell $(ROCM_LLVM)/../bin/rocm_agent_enumerator | grep -m 1 -E gfx[^0]{1})
+ROCM_GPU       ?= $(INSTALLED_GPU)
+endif
+
+ifeq ($(TARGETS),)
+TARGETS =--offload-arch=$(ROCM_GPU):xnack+
+endif
+
+CC       =$(ROCM_LLVM)/bin/clang++
+
+VERS = $(shell $(ROCM_LLVM)/bin/clang --version | grep -oP '(?<=clang version )[0-9.]+')
+ROCM = $(shell $(ROCM_LLVM)/bin/clang --version | grep -oP 'AMD clang')
+
+ifeq ($(ROCM),)
+  ROCMHIP ?= $(ROCM_LLVM)
+else
+  ROCMHIP ?= $(ROCM_LLVM)/..
+endif
+
+ifeq ($(shell expr $(VERS) \>= 12.0), 1)
+        RPTH = -Wl,-rpath,$(ROCMHIP)/lib -Wl,-rpath,$(ROCM_LLVM)/lib/clang/$(VERS)/lib/linux
+endif
+
+#  compiler automatically adds "libdevice/<target>/ to -L opts
+LFLAGS   =-L$(ROCMHIP)/lib $(RPTH)
+ASAN_FLAGS = -fsanitize=address -shared-libsan
+CFLAGS = -O0 -g -std=c++11 -fopenmp 
+
+# ----- Demo compile and link in one step, no object code saved
+$(TESTNAME): $(TESTNAME).$(FILETYPE)
+	$(CC) $(CFLAGS) $(TARGETS) $(ASAN_FLAGS) $(LFLAGS) $^ -o $@
+
+run: $(TESTNAME)
+	HSA_XNACK=1 ./$(TESTNAME)
+
+#  ----   Demo compile and link in two steps, object saved
+$(TESTNAME).o: $(TESTNAME).$(FILETYPE)
+	$(CC) -c $(CFLAGS) $(TARGETS) $(ASAN_FLAGS) $^
+
+obin:	$(TESTNAME).o
+	$(CC) $(ASAN_FLAGS) $(LFLAGS) $^ -o obin
+
+run_obin: obin
+	./obin
+
+#  ----   Demo compile to intermediates LLVMIR or assembly
+$(TESTNAME).ll: $(TESTNAME).$(FILETYPE)
+	$(CC) -c -S -emit-llvm $(CFLAGS) $(TARGETS) $(ASAN_FLAGS) $^
+
+$(TESTNAME).s: $(TESTNAME).$(FILETYPE)
+	$(CC)  -c -S $(CFLAGS) $(TARGETS) $(ASAN_FLAGS) $^
+
+help:
+	@echo
+	@echo "Makefile Help:"
+	@echo "  Source:		$(TESTNAME).$(FILETYPE)"
+	@echo "  Compiler: 		$(CC)"
+	@echo "  Compiler flags: 	$(CFLAGS)"
+	@echo
+	@echo "Avalable Targets:"
+	@echo "  make			// build binary $(TESTNAME) with ASan flags"
+	@echo "  make run		// run $(TESTNAME) with ASan flags"
+	@echo "  make $(TESTNAME).o	// compile, be, & assemble : -c"
+	@echo "  make obin         	// link step only"
+	@echo "  make run_obin     	// run obin "
+	@echo "  make $(TESTNAME).s	// compile & backend steps : -c -S"
+	@echo "  make $(TESTNAME).ll	// compile step only       : -c -S -emit-llvm"
+	@echo "  make clean		// cleanup files"
+	@echo "  make help 		// this help"
+	@echo
+	@echo "Environment Variables:"
+	@echo "  ROCM_LLVM  default: /opt/rocm/llvm     value: $(ROCM_LLVM)"
+	@echo "  ROCM_GPU   default: $(INSTALLED_GPU)             value: $(ROCM_GPU)"
+	@echo "  TARGETS    default: --offload-arch=$(ROCM_GPU)"
+	@echo "              value: $(TARGETS)"
+	@echo
+	@echo "Link Flags:"
+	@echo "  Link flags: $(LFLAGS)"
+	@echo
+
+# Cleanup anything this makefile can create
+clean:
+	@[ -f ./$(TESTNAME) ] && rm ./$(TESTNAME) ; true
+	@[ -f ./obin ] && rm ./obin ; true
+	@[ -f ./$(TESTNAME).ll ] && rm *.ll ; true
+	@[ -f ./$(TESTNAME).o ] && rm $(TESTNAME).o ; true
+	@[ -f ./$(TESTNAME).s ] && rm *.s ; true
diff --git a/examples/tools/asan/global_buffer_overflow/openmp/vecadd-GBO.cpp b/examples/tools/asan/global_buffer_overflow/openmp/vecadd-GBO.cpp
@@ -0,0 +1,65 @@
+// MIT License
+//
+// Copyright (c) 2017 Advanced Micro Devices, Inc. All Rights Reserved.
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use, copy,
+// modify, merge, publish, distribute, sublicense, and/or sell copies
+// of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS
+// BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN
+// ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
+// CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include <cstdio>
+#include <cstdlib>
+#include <omp.h>
+
+#define N 100
+
+#pragma omp declare target
+int A[N], B[N], C[N];
+#pragma omp end declare target
+
+void Print(int arr[], int n) {
+  for (int i = 0; i < n; i++) {
+    printf("\n%d", arr[i]);
+  }
+}
+
+int main(int argc, char *argv[]) {
+
+  for (int i = 0; i < N; i++) {
+    A[i] = 2 * (i + 1);
+    B[i] = 3 * (i + 1);
+  }
+
+#pragma omp target data map(to : A [0:N], B [0:N]) map(from : C [0:N])
+  {
+#pragma omp target update to(A, B)
+
+#pragma omp target parallel for
+    for (int i = 0; i < N; i++) {
+      C[i] = A[i + 10] + B[i];
+    }
+
+#pragma omp target update from(C)
+  }
+
+  Print(C, N);
+
+  printf("\n");
+
+  return 0;
+}