Use new builtins

Change-Id: I201fadf7c477713cdaf21207eed8e1cac7746134

ockl/src/toas.cl

@@ -1,39 +1,39 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
 
-#include "irif.h"
-#include "device_amd_hsa.h"
 #include "ockl.h"
 
-__attribute__((always_inline, const)) bool
+__attribute__((const, target("flat-address-space"))) bool
 OCKL_MANGLE_T(is_local,addr)(const void *a)
 {
-    __constant amd_queue_t *q = (__constant amd_queue_t *)__builtin_amdgcn_queue_ptr();
-    uint u = (uint)((ulong)a >> 32);
-    return u == q->group_segment_aperture_base_hi;
+    return __builtin_amdgcn_is_shared(a);
 }
 
-__attribute__((always_inline, const)) bool
+__attribute__((const, target("flat-address-space"))) bool
 OCKL_MANGLE_T(is_private,addr)(const void *a)
 {
-    __constant amd_queue_t *q = (__constant amd_queue_t *)__builtin_amdgcn_queue_ptr();
-    uint u = (uint)((ulong)a >> 32);
-    return u == q->private_segment_aperture_base_hi;
+    return __builtin_amdgcn_is_private(a);
 }
 
-__attribute__((always_inline, const)) __global void *
+__attribute__((const)) __global void *
 OCKL_MANGLE_T(to,global)(void *a)
 {
     __global void *ga = (__global void *)((ulong)a);
     return OCKL_MANGLE_T(is_local,addr)(a) | OCKL_MANGLE_T(is_private,addr)(a) ?  (__global void *)0UL : ga;
 }
 
-__attribute__((always_inline, const)) __local void *
+__attribute__((const)) __local void *
 OCKL_MANGLE_T(to,local)(void *a)
 {
     uint u = (uint)((ulong)a);
     return OCKL_MANGLE_T(is_local,addr)(a) ? (__local void *)u : (__local void *)0;
 }
 
-__attribute__((always_inline, const)) __private void *
+__attribute__((const)) __private void *
 OCKL_MANGLE_T(to,private)(void *a)
 {
     uint u = (uint)((ulong)a);