diff options
| author | Peter Collingbourne <peter@pcc.me.uk> | 2012-01-08 22:09:58 +0000 |
|---|---|---|
| committer | Peter Collingbourne <peter@pcc.me.uk> | 2012-01-08 22:09:58 +0000 |
| commit | d5395fbf03dc168d00bc8622e86fefbac9a3da2a (patch) | |
| tree | 2c24697a31ef5ef5445a8ad143aa34f4ef42db9d /libclc/ptx-nvidiacl | |
| parent | 11faafe7dce2bd8befa3e2444c2292edcd9b2d50 (diff) | |
| download | llvm-d5395fbf03dc168d00bc8622e86fefbac9a3da2a.tar.gz | |
Initial commit.
llvm-svn: 147756
Diffstat (limited to 'libclc/ptx-nvidiacl')
8 files changed, 54 insertions, 0 deletions
diff --git a/libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h b/libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h new file mode 100644 index 000000000000..cd9f3276af62 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h @@ -0,0 +1,6 @@ +_CLC_INLINE void barrier(cl_mem_fence_flags flags) { + if (flags & CLK_LOCAL_MEM_FENCE) { + __builtin_ptx_bar_sync(0); + } +} + diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h new file mode 100644 index 000000000000..026d2fea31f8 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_global_id(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_ctaid_x()*__builtin_ptx_read_ntid_x()+__builtin_ptx_read_tid_x(); + case 1: return __builtin_ptx_read_ctaid_y()*__builtin_ptx_read_ntid_y()+__builtin_ptx_read_tid_y(); + case 2: return __builtin_ptx_read_ctaid_z()*__builtin_ptx_read_ntid_z()+__builtin_ptx_read_tid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h new file mode 100644 index 000000000000..5cd4222d5ae7 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_global_size(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_nctaid_x()*__builtin_ptx_read_ntid_x(); + case 1: return __builtin_ptx_read_nctaid_y()*__builtin_ptx_read_ntid_y(); + case 2: return __builtin_ptx_read_nctaid_z()*__builtin_ptx_read_ntid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h new file mode 100644 index 000000000000..18b1bd4db079 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_group_id(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_ctaid_x(); + case 1: return __builtin_ptx_read_ctaid_y(); + case 2: return __builtin_ptx_read_ctaid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h new file mode 100644 index 000000000000..1b8c776badf6 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_local_id(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_tid_x(); + case 1: return __builtin_ptx_read_tid_y(); + case 2: return __builtin_ptx_read_tid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h new file mode 100644 index 000000000000..cbc1f6ee2893 --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_local_size(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_ntid_x(); + case 1: return __builtin_ptx_read_ntid_y(); + case 2: return __builtin_ptx_read_ntid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h new file mode 100644 index 000000000000..36ee849c613f --- /dev/null +++ b/libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h @@ -0,0 +1,8 @@ +_CLC_INLINE size_t get_num_groups(uint dim) { + switch (dim) { + case 0: return __builtin_ptx_read_nctaid_x(); + case 1: return __builtin_ptx_read_nctaid_y(); + case 2: return __builtin_ptx_read_nctaid_z(); + default: return 0; + } +} diff --git a/libclc/ptx-nvidiacl/lib/SOURCES b/libclc/ptx-nvidiacl/lib/SOURCES new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/libclc/ptx-nvidiacl/lib/SOURCES |