Adjust address to point at or within the call

Change-Id: Id142d6c0c8e07622a34dc03a86a9161a6fbee199

Author: b-sumner

amd/device-libs/asanrtl/inc/asan_util.h

@@ -19,7 +19,8 @@ typedef unsigned long u64;
 
 #define SHADOW_GRANULARITY (1ULL << ASAN_SHADOW)
 
-#define GET_CALLER_PC() (uptr) __builtin_return_address(0)
+#define CALL_BYTES 4
+#define GET_CALLER_PC() ((uptr)__builtin_return_address(0) - CALL_BYTES)
 
 #define WORKGROUP_ID(dim) __builtin_amdgcn_workgroup_id_##dim()
 

amd/device-libs/asanrtl/src/dm.cl

@@ -333,6 +333,8 @@ __asan_free_impl(ulong aa, ulong pc)
     if (!aa)
         return;
 
+    pc -= CALL_BYTES;
+
     uptr sa = MEM_TO_SHADOW(aa);
     s8 sb = *(__global s8*) sa;
     if (sb != 0 && ((s8)(aa & (SHADOW_GRANULARITY-1)) >= sb)) {
@@ -613,6 +615,8 @@ NO_SANITIZE_ADDR
 ulong
 __asan_malloc_impl(ulong sz, ulong pc)
 {
+    pc -= CALL_BYTES;
+
     if (sz > SLAB_THRESHOLD)
         return non_slab_malloc(sz, pc);
     else