diff options
| author | Andrzej Janik <[email protected]> | 2021-09-17 20:53:44 +0000 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2021-09-17 20:53:44 +0000 |
| commit | d5a4b068dd9bf72a0e1b6448583ebad609ed72c1 (patch) | |
| tree | 988aaa9772d78a0963bb54a96f93f95de15877b5 /ptx/lib | |
| parent | 6ef19d65010164a7cc8408663eb189b64f44d26a (diff) | |
| download | ZLUDA-d5a4b068dd9bf72a0e1b6448583ebad609ed72c1.tar.gz ZLUDA-d5a4b068dd9bf72a0e1b6448583ebad609ed72c1.zip | |
Redo handling of sregs
Diffstat (limited to 'ptx/lib')
| -rw-r--r-- | ptx/lib/zluda_ptx_impl.bc | bin | 31940 -> 33884 bytes | |||
| -rw-r--r-- | ptx/lib/zluda_ptx_impl.cl | 16 |
2 files changed, 16 insertions, 0 deletions
diff --git a/ptx/lib/zluda_ptx_impl.bc b/ptx/lib/zluda_ptx_impl.bc Binary files differindex 7aa12c8..e2f956d 100644 --- a/ptx/lib/zluda_ptx_impl.bc +++ b/ptx/lib/zluda_ptx_impl.bc diff --git a/ptx/lib/zluda_ptx_impl.cl b/ptx/lib/zluda_ptx_impl.cl index d439795..0870fb5 100644 --- a/ptx/lib/zluda_ptx_impl.cl +++ b/ptx/lib/zluda_ptx_impl.cl @@ -297,6 +297,22 @@ atomic_add(atom_acq_rel_sys_shared_add_f64, memory_order_acq_rel, memory_order_a return (uint)__builtin_amdgcn_uicmp(1, 0, 33);
}
+ uint FUNC(sreg_tid)(uchar dim) {
+ return (uint)get_local_id(dim);
+ }
+
+ uint FUNC(sreg_ntid)(uchar dim) {
+ return (uint)get_local_size(dim);
+ }
+
+ uint FUNC(sreg_ctaid)(uchar dim) {
+ return (uint)get_group_id(dim);
+ }
+
+ uint FUNC(sreg_nctaid)(uchar dim) {
+ return (uint)get_num_groups(dim);
+ }
+
uint FUNC(sreg_clock)() {
return (uint)__builtin_amdgcn_s_memtime();
}
|