merge amd-stg-open into amd-mainline-open

Change-Id: I946330c7455ee86a77183cfef8ed64479568963a

asanrtl/inc/asan_util.h

@@ -0,0 +1,60 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#pragma once
+#include "ockl.h"
+
+typedef ulong uptr;
+typedef unsigned char u8;
+typedef signed char s8;
+typedef unsigned short u16;
+typedef short s16;
+typedef unsigned long u64;
+
+#define ASAN_SHADOW 3
+
+#define SHADOW_GRANULARITY (1ULL << ASAN_SHADOW)
+
+#define GET_CALLER_PC() (uptr) __builtin_return_address(0)
+
+#define WORKGROUP_ID(dim) __builtin_amdgcn_workgroup_id_##dim()
+
+#define OPT_NONE __attribute__((optnone))
+
+#define NO_SANITIZE_ADDR __attribute__((no_sanitize("address")))
+
+#define REPORT_IMPL(caller_pc, addr, is_write, size, no_abort)                 \
+    uptr read = is_write;                                                      \
+    if (no_abort)                                                              \
+        read |= 0xFFFFFFFF00000000;                                            \
+                                                                               \
+    __ockl_sanitizer_report(addr, caller_pc, WORKGROUP_ID(x), WORKGROUP_ID(y), \
+                            WORKGROUP_ID(z), __ockl_get_local_linear_id(),     \
+                            read, size);
+
+NO_SANITIZE_ADDR
+static bool
+is_aligned_by_granularity(uptr addr)
+{
+    return (addr & (SHADOW_GRANULARITY - 1)) == 0;
+}
+
+// round up size to the nearest multiple of boundary.
+NO_SANITIZE_ADDR
+static uptr
+round_upto(uptr size, uptr boundary)
+{
+    return (size + boundary - 1) & ~(boundary - 1);
+}
+
+// round down size to the nearest multiple of boundary.
+NO_SANITIZE_ADDR
+static uptr
+round_downto(uptr size, uptr boundary)
+{
+    return size & ~(boundary - 1);
+}

asanrtl/inc/globals.h

@@ -6,7 +6,7 @@
  *===------------------------------------------------------------------------*/
 
 #pragma once
-#include "shadow_mapping.h"
+#include "asan_util.h"
 
 // The strucutures semantics and layout must match the host instrumented
 // global variable as defined in

asanrtl/inc/shadow_mapping.h

@@ -6,22 +6,16 @@
  *===------------------------------------------------------------------------*/
 
 #pragma once
-#include "asan_type_decls.h"
-
-#define ASAN_SHADOW 3
-
-#define SHADOW_GRANULARITY (1ULL << ASAN_SHADOW)
+#include "asan_util.h"
 
 //offset from llvm/compiler-rt/lib/asan/asan_mapping.h
 static const u64 kh_Linux64bit_ShadowOffset =
     0x7FFFFFFF & (~0xFFFULL << ASAN_SHADOW);
 
 #define MEM_TO_SHADOW(mem_addr) (((mem_addr) >> ASAN_SHADOW) + kh_Linux64bit_ShadowOffset)
 
-#define NO_SANITIZE_ADDR __attribute__((no_sanitize("address")))
-
-//address are atleast SHADOW_GRANULARITY aligned
-//true, when given byte is accessible false otherwise
+// Addresses are atleast SHADOW_GRANULARITY aligned.
+// True, when given byte is accessible false otherwise.
 NO_SANITIZE_ADDR
 static bool
 is_address_poisoned(uptr addr)
@@ -35,21 +29,6 @@ is_address_poisoned(uptr addr)
     return false;
 }
 
-//check all application bytes in [beg,beg+size) range are accessible
 NO_SANITIZE_ADDR
-static bool
-is_region_poisoned(uptr beg, uptr size)
-{
-    uptr end = beg + size - 1;
-    // Fast path - check first and last application bytes
-    if (is_address_poisoned(beg) ||
-        is_address_poisoned(end))
-    return true;
-
-    // check all inner bytes
-    for (uptr addr = beg+1; addr < end; addr++){
-       if (is_address_poisoned(addr))
-        return true;
-    }
-    return false;
-}
+uptr
+__asan_region_is_poisoned(uptr beg, uptr size);

asanrtl/src/globals.cl

@@ -5,24 +5,9 @@
  * License. See LICENSE.TXT for details.
  *===------------------------------------------------------------------------*/
 
+#include "asan_util.h"
 #include "globals.h"
-
-static bool
-is_aligned_by_granularity(uptr addr) {
-    return (addr & (SHADOW_GRANULARITY - 1)) == 0;
-}
-
-// round up size to the nearest multiple of boundary.
-static uptr
-round_upto(uptr size, uptr boundary) {
-    return (size + boundary - 1) & ~(boundary - 1);
-}
-
-// round down size to the nearest multiple of boundary.
-static uptr
-round_downto(uptr size, uptr boundary) {
-    return size & ~(boundary - 1);
-}
+#include "shadow_mapping.h"
 
 // fill shadow bytes of range [aligned_beg, aligned_beg+aligned_size)
 // with value.

asanrtl/src/memintrinsics.cl

@@ -0,0 +1,61 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "asan_util.h"
+#include "shadow_mapping.h"
+
+OPT_NONE
+NO_SANITIZE_ADDR
+static void
+check_memory_range_accessible(void* dest, const void* src,
+                    uptr size, uptr pc) {
+    if (size == 0)
+      return;
+    uptr invalid_addr = 0;
+    uptr src_addr = (uptr)src;
+    invalid_addr = __asan_region_is_poisoned(src_addr, size);
+    if (invalid_addr) {
+      REPORT_IMPL(pc, invalid_addr, false, size, false)
+    }
+    uptr dest_addr = (uptr)dest;
+    invalid_addr = __asan_region_is_poisoned(dest_addr, size);
+    if (invalid_addr) {
+      REPORT_IMPL(pc, invalid_addr, true, size, false)
+    }
+}
+
+OPT_NONE
+NO_SANITIZE_ADDR
+void*
+__asan_memcpy(void* to, const void* from, uptr size) {
+    uptr pc = GET_CALLER_PC();
+    check_memory_range_accessible(to, from, size, pc);
+    return __builtin_memcpy(to, from, size);
+}
+
+OPT_NONE
+NO_SANITIZE_ADDR
+void*
+__asan_memmove(void* to, const void* from, uptr size) {
+    uptr pc = GET_CALLER_PC();
+    check_memory_range_accessible(to, from, size, pc);
+    return __builtin_memmove(to, from, size);
+}
+
+OPT_NONE
+NO_SANITIZE_ADDR
+void*
+__asan_memset(void* s, int c, uptr n) {
+    uptr pc = GET_CALLER_PC();
+    uptr src_addr = (uptr)s;
+    uptr invalid_addr = 0;
+    invalid_addr = __asan_region_is_poisoned(src_addr, n);
+    if (invalid_addr) {
+      REPORT_IMPL(pc, invalid_addr, true, n, false)
+    }
+    return __builtin_memset(s, c, n);
+}

asanrtl/src/preserve.cl

@@ -54,6 +54,7 @@ extern void __asan_store_n (uptr addr, uptr size);
 extern void __asan_store_n_noabort (uptr addr, uptr size);
 extern void __asan_load_n (uptr addr, uptr size);
 extern void __asan_load_n_noabort (uptr addr, uptr size);
+extern uptr __asan_region_is_poisoned(uptr beg, uptr size);
 extern void* __asan_memmove(void* to, void* from, uptr size);
 extern void* __asan_memcpy(void* to, void* from, uptr size);
 extern void* __asan_memset(void* s, int c, uptr n);
@@ -122,6 +123,7 @@ __amdgpu_device_library_preserve_asan_functions(void)
     __asan_store_n_noabort(0, 0);
     __asan_load_n(0, 0);
     __asan_load_n_noabort(0, 0);
+    __asan_region_is_poisoned(0, 0);
     (void)__asan_memmove((void*)0, (void*)0, 0);
     (void)__asan_memcpy((void*)0, (void*)0, 0);
     (void)__asan_memset((void*)0, 0, 0);

asanrtl/src/report.cl

@@ -5,27 +5,9 @@
  * License. See LICENSE.TXT for details.
  *===------------------------------------------------------------------------*/
 
-#include "ockl.h"
-#include "asan_type_decls.h"
+#include "asan_util.h"
 #include "shadow_mapping.h"
 
-#define GET_CALLER_PC() \
-    (uptr) __builtin_return_address(0)
-
-#define WORKGROUP_ID(dim) \
-    __builtin_amdgcn_workgroup_id_ ##dim()
-
-#define OPT_NONE __attribute__((optnone))
-
-#define REPORT_IMPL(caller_pc, addr, is_write, size, no_abort)                          \
-    uptr read      = is_write;                                                          \
-    if (no_abort)                                                                       \
-        read       |= 0xFFFFFFFF00000000;                                               \
-                                                                                        \
-    __ockl_sanitizer_report(addr, caller_pc, WORKGROUP_ID(x), WORKGROUP_ID(y),          \
-                          WORKGROUP_ID(z), __ockl_get_local_linear_id(), read, size);   \
-
-
 #define ASAN_REPORT_ERROR(type, size, is_write)                                  \
 OPT_NONE                                                                         \
 void __asan_report_ ## type ## size(uptr addr) {                                 \
@@ -108,14 +90,14 @@ ASAN_ERROR(store, 16,1)
 OPT_NONE NO_SANITIZE_ADDR                                    \
 void __asan_ ## type ## _n(uptr addr, uptr size) {           \
     uptr caller_pc = GET_CALLER_PC();                        \
-    if (is_region_poisoned(addr, size)) {                    \
+    if (__asan_region_is_poisoned(addr, size)) {             \
         REPORT_IMPL(caller_pc, addr, is_write, size, false)  \
     }                                                        \
 }                                                            \
 OPT_NONE NO_SANITIZE_ADDR                                    \
 void __asan_ ## type ## _n_noabort(uptr addr, uptr size) {   \
     uptr caller_pc = GET_CALLER_PC();                        \
-    if (is_region_poisoned(addr, size)) {                    \
+    if (__asan_region_is_poisoned(addr, size)) {             \
         REPORT_IMPL(caller_pc, addr, is_write, size, true)   \
     }                                                        \
 }                                                            \

asanrtl/src/shadow_mapping.cl

@@ -0,0 +1,47 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "shadow_mapping.h"
+
+NO_SANITIZE_ADDR
+static uptr
+range_check(uptr beg, uptr end) {
+    uptr aligned_beg = round_downto(beg, SHADOW_GRANULARITY);
+    uptr aligned_end = round_downto(end, SHADOW_GRANULARITY);
+    uptr shadow_beg  = MEM_TO_SHADOW(aligned_beg);
+    uptr shadow_end  = MEM_TO_SHADOW(aligned_end);
+    uptr nbytes      = (shadow_end - shadow_beg)+1;
+    uptr shadow_byte_count = 0;
+    while (shadow_beg <= shadow_end) {
+      s8 shadow_value = *(__global s8 *)shadow_beg;
+      if (shadow_value)
+        break;
+      shadow_byte_count++;
+      shadow_beg++;
+    }
+    if (shadow_byte_count == nbytes)
+      return 0;
+    uptr start_addr = round_downto(beg + (shadow_byte_count*SHADOW_GRANULARITY), SHADOW_GRANULARITY);
+    return start_addr;
+}
+
+//check all application bytes in [beg,beg+size) range are accessible
+NO_SANITIZE_ADDR
+uptr
+__asan_region_is_poisoned(uptr beg, uptr size)
+{
+    uptr end  = beg + size - 1;
+    uptr start_addr = range_check(beg, end);
+    if (start_addr != 0) {
+      // loop through the range to find accessible address.
+      for (uptr addr = start_addr; addr <= end; ++addr) {
+        if (is_address_poisoned(addr))
+          return addr;
+      }
+    }
+    return 0;
+}

asanrtl/src/stubs.cl

@@ -7,13 +7,6 @@
 
 typedef ulong uptr;
 
-
-void* __asan_memmove(void* to, void* from, uptr size) { return to; }
-
-void* __asan_memcpy(void* to, void* from, uptr size) { return to; }
-
-void* __asan_memset(void* s, int c, uptr n) { return s; }
-
 void __asan_handle_no_return(void) {}
 
 void __sanitizer_ptr_cmp(uptr a, uptr b) {}

ockl/src/services.cl

@@ -401,7 +401,7 @@ __ockl_sanitizer_report(ulong addr, ulong pc, ulong wgidx, ulong wgidy,
 
 /*===---  DEVMEM  ----------------------------------------------------------*/
 
-ulong
+WEAK_ATTR ulong
 __ockl_devmem_request(ulong addr, ulong size)
 {
     long2 result = __ockl_hostcall_preview(SERVICE_DEVMEM, addr, size, 0, 0, 0, 0, 0, 0);