diff options
author | Peter Collingbourne <peter@pcc.me.uk> | 2012-08-05 22:25:37 +0000 |
---|---|---|
committer | Peter Collingbourne <peter@pcc.me.uk> | 2012-08-05 22:25:37 +0000 |
commit | a385c5341329da54e4834aa0af5e382617114d3f (patch) | |
tree | 48b74090743ab20ecd0bbf04e65b8cec784885a5 /libclc/ptx-nvidiacl | |
parent | 1e373f07af5113529cb03f4699971b492cb508ed (diff) | |
download | bcm5719-llvm-a385c5341329da54e4834aa0af5e382617114d3f.tar.gz bcm5719-llvm-a385c5341329da54e4834aa0af5e382617114d3f.zip |
PTX: move implementations of work-item and synchronisation functions
to lib, and add header files in generic. Incorporates a patch by
Tom Stellard!
llvm-svn: 161313
Diffstat (limited to 'libclc/ptx-nvidiacl')
-rw-r--r-- | libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h | 8 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h | 8 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/SOURCES | 4 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/synchronization/barrier.cl (renamed from libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h) | 4 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl (renamed from libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h) | 4 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl (renamed from libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h) | 4 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl (renamed from libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h) | 4 | ||||
-rw-r--r-- | libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl (renamed from libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h) | 4 |
8 files changed, 19 insertions, 21 deletions
diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h b/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h deleted file mode 100644 index 026d2fea31f..00000000000 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h +++ /dev/null @@ -1,8 +0,0 @@ -_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 deleted file mode 100644 index 5cd4222d5ae..00000000000 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h +++ /dev/null @@ -1,8 +0,0 @@ -_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/lib/SOURCES b/libclc/ptx-nvidiacl/lib/SOURCES index e69de29bb2d..1a96a1ab873 100644 --- a/libclc/ptx-nvidiacl/lib/SOURCES +++ b/libclc/ptx-nvidiacl/lib/SOURCES @@ -0,0 +1,4 @@ +workitem/get_group_id.cl +workitem/get_local_id.cl +workitem/get_local_size.cl +workitem/get_num_groups.cl diff --git a/libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h b/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl index cd9f3276af6..fb36c2612be 100644 --- a/libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h +++ b/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl @@ -1,4 +1,6 @@ -_CLC_INLINE void barrier(cl_mem_fence_flags flags) { +#include <clc/clc.h> + +_CLC_DEF 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_group_id.h b/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl index 18b1bd4db07..2b35b4eaaa9 100644 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h +++ b/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl @@ -1,4 +1,6 @@ -_CLC_INLINE size_t get_group_id(uint dim) { +#include <clc/clc.h> + +_CLC_DEF 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(); diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h b/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl index 1b8c776badf..f0cfdc005fe 100644 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h +++ b/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl @@ -1,4 +1,6 @@ -_CLC_INLINE size_t get_local_id(uint dim) { +#include <clc/clc.h> + +_CLC_DEF 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(); diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h b/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl index cbc1f6ee289..c3f542595de 100644 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h +++ b/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl @@ -1,4 +1,6 @@ -_CLC_INLINE size_t get_local_size(uint dim) { +#include <clc/clc.h> + +_CLC_DEF 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(); diff --git a/libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h b/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl index 36ee849c613..90bdc2e41d2 100644 --- a/libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h +++ b/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl @@ -1,4 +1,6 @@ -_CLC_INLINE size_t get_num_groups(uint dim) { +#include <clc/clc.h> + +_CLC_DEF 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(); |