aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx/lib/zluda_ptx_impl.cl
diff options
context:
space:
mode:
authorAndrzej Janik <[email protected]>2021-09-17 16:24:25 +0000
committerAndrzej Janik <[email protected]>2021-09-17 16:24:25 +0000
commit5b2352723fb251b64317737167b609a0a11651a6 (patch)
tree35562f2fdf7b39b56ebe0bb74113b5f1ff2d8d88 /ptx/lib/zluda_ptx_impl.cl
parentc37223fe673f2f45a533e338b74ae9325748588a (diff)
downloadZLUDA-5b2352723fb251b64317737167b609a0a11651a6.tar.gz
ZLUDA-5b2352723fb251b64317737167b609a0a11651a6.zip
Implement function pointers and activemask
Diffstat (limited to 'ptx/lib/zluda_ptx_impl.cl')
-rw-r--r--ptx/lib/zluda_ptx_impl.cl5
1 files changed, 5 insertions, 0 deletions
diff --git a/ptx/lib/zluda_ptx_impl.cl b/ptx/lib/zluda_ptx_impl.cl
index 9171ef9..aca9327 100644
--- a/ptx/lib/zluda_ptx_impl.cl
+++ b/ptx/lib/zluda_ptx_impl.cl
@@ -291,6 +291,11 @@ atomic_add(atom_acq_rel_sys_shared_add_f64, memory_order_acq_rel, memory_order_a
ulong FUNC(brev_b64)(ulong base) {
return __llvm_bitreverse_i64(base);
}
+
+ // Taken from __ballot definition in hipamd/include/hip/amd_detail/amd_device_functions.h
+ uint FUNC(activemask)() {
+ return (uint)__builtin_amdgcn_uicmp(1, 0, 33);
+ }
#endif
void FUNC(__assertfail)(