From a385c5341329da54e4834aa0af5e382617114d3f Mon Sep 17 00:00:00 2001 From: Peter Collingbourne Date: Sun, 5 Aug 2012 22:25:37 +0000 Subject: [PATCH] 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 --- libclc/generic/include/clc/synchronization/barrier.h | 1 + libclc/generic/include/clc/workitem/get_global_id.h | 1 + libclc/generic/include/clc/workitem/get_global_size.h | 1 + libclc/generic/include/clc/workitem/get_group_id.h | 1 + libclc/generic/include/clc/workitem/get_local_id.h | 1 + libclc/generic/include/clc/workitem/get_local_size.h | 1 + libclc/generic/include/clc/workitem/get_num_groups.h | 1 + libclc/generic/lib/SOURCES | 2 ++ libclc/generic/lib/workitem/get_global_id.cl | 5 +++++ libclc/generic/lib/workitem/get_global_size.cl | 5 +++++ libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h | 8 -------- libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h | 8 -------- libclc/ptx-nvidiacl/lib/SOURCES | 4 ++++ .../synchronization/barrier.h => lib/synchronization/barrier.cl} | 4 +++- .../clc/workitem/get_group_id.h => lib/workitem/get_group_id.cl} | 4 +++- .../clc/workitem/get_local_id.h => lib/workitem/get_local_id.cl} | 4 +++- .../workitem/get_local_size.h => lib/workitem/get_local_size.cl} | 4 +++- .../workitem/get_num_groups.h => lib/workitem/get_num_groups.cl} | 4 +++- 18 files changed, 38 insertions(+), 21 deletions(-) create mode 100644 libclc/generic/include/clc/synchronization/barrier.h create mode 100644 libclc/generic/include/clc/workitem/get_global_id.h create mode 100644 libclc/generic/include/clc/workitem/get_global_size.h create mode 100644 libclc/generic/include/clc/workitem/get_group_id.h create mode 100644 libclc/generic/include/clc/workitem/get_local_id.h create mode 100644 libclc/generic/include/clc/workitem/get_local_size.h create mode 100644 libclc/generic/include/clc/workitem/get_num_groups.h create mode 100644 libclc/generic/lib/workitem/get_global_id.cl create mode 100644 libclc/generic/lib/workitem/get_global_size.cl delete mode 100644 libclc/ptx-nvidiacl/include/clc/workitem/get_global_id.h delete mode 100644 libclc/ptx-nvidiacl/include/clc/workitem/get_global_size.h rename libclc/ptx-nvidiacl/{include/clc/synchronization/barrier.h => lib/synchronization/barrier.cl} (51%) rename libclc/ptx-nvidiacl/{include/clc/workitem/get_group_id.h => lib/workitem/get_group_id.cl} (74%) rename libclc/ptx-nvidiacl/{include/clc/workitem/get_local_id.h => lib/workitem/get_local_id.cl} (74%) rename libclc/ptx-nvidiacl/{include/clc/workitem/get_local_size.h => lib/workitem/get_local_size.cl} (74%) rename libclc/ptx-nvidiacl/{include/clc/workitem/get_num_groups.h => lib/workitem/get_num_groups.cl} (74%) diff --git a/libclc/generic/include/clc/synchronization/barrier.h b/libclc/generic/include/clc/synchronization/barrier.h new file mode 100644 index 0000000..7167a3d --- /dev/null +++ b/libclc/generic/include/clc/synchronization/barrier.h @@ -0,0 +1 @@ +_CLC_DECL void barrier(cl_mem_fence_flags flags); diff --git a/libclc/generic/include/clc/workitem/get_global_id.h b/libclc/generic/include/clc/workitem/get_global_id.h new file mode 100644 index 0000000..92759f1 --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_global_id.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_global_id(uint dim); diff --git a/libclc/generic/include/clc/workitem/get_global_size.h b/libclc/generic/include/clc/workitem/get_global_size.h new file mode 100644 index 0000000..2f83705 --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_global_size.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_global_size(uint dim); diff --git a/libclc/generic/include/clc/workitem/get_group_id.h b/libclc/generic/include/clc/workitem/get_group_id.h new file mode 100644 index 0000000..346c82c --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_group_id.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_group_id(uint dim); diff --git a/libclc/generic/include/clc/workitem/get_local_id.h b/libclc/generic/include/clc/workitem/get_local_id.h new file mode 100644 index 0000000..169aeed --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_local_id.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_local_id(uint dim); diff --git a/libclc/generic/include/clc/workitem/get_local_size.h b/libclc/generic/include/clc/workitem/get_local_size.h new file mode 100644 index 0000000..040ec58 --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_local_size.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_local_size(uint dim); diff --git a/libclc/generic/include/clc/workitem/get_num_groups.h b/libclc/generic/include/clc/workitem/get_num_groups.h new file mode 100644 index 0000000..e555c7e --- /dev/null +++ b/libclc/generic/include/clc/workitem/get_num_groups.h @@ -0,0 +1 @@ +_CLC_DECL size_t get_num_groups(uint dim); diff --git a/libclc/generic/lib/SOURCES b/libclc/generic/lib/SOURCES index 344c865..1d56c40 100644 --- a/libclc/generic/lib/SOURCES +++ b/libclc/generic/lib/SOURCES @@ -12,3 +12,5 @@ integer/sub_sat.ll integer/sub_sat_impl.ll math/hypot.cl math/mad.cl +workitem/get_global_id.cl +workitem/get_global_size.cl diff --git a/libclc/generic/lib/workitem/get_global_id.cl b/libclc/generic/lib/workitem/get_global_id.cl new file mode 100644 index 0000000..fdd83d2 --- /dev/null +++ b/libclc/generic/lib/workitem/get_global_id.cl @@ -0,0 +1,5 @@ +#include + +_CLC_DEF size_t get_global_id(uint dim) { + return get_group_id(dim)*get_local_size(dim) + get_local_id(dim); +} diff --git a/libclc/generic/lib/workitem/get_global_size.cl b/libclc/generic/lib/workitem/get_global_size.cl new file mode 100644 index 0000000..5ae649e --- /dev/null +++ b/libclc/generic/lib/workitem/get_global_size.cl @@ -0,0 +1,5 @@ +#include + +_CLC_DEF size_t get_global_size(uint dim) { + return get_num_groups(dim)*get_local_size(dim); +} 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 026d2fe..0000000 --- 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 5cd4222..0000000 --- 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 e69de29..1a96a1a 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 similarity index 51% rename from libclc/ptx-nvidiacl/include/clc/synchronization/barrier.h rename to libclc/ptx-nvidiacl/lib/synchronization/barrier.cl index cd9f327..fb36c26 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_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 similarity index 74% rename from libclc/ptx-nvidiacl/include/clc/workitem/get_group_id.h rename to libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl index 18b1bd4..2b35b4e 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_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 similarity index 74% rename from libclc/ptx-nvidiacl/include/clc/workitem/get_local_id.h rename to libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl index 1b8c776..f0cfdc0 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_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 similarity index 74% rename from libclc/ptx-nvidiacl/include/clc/workitem/get_local_size.h rename to libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl index cbc1f6e..c3f5425 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_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 similarity index 74% rename from libclc/ptx-nvidiacl/include/clc/workitem/get_num_groups.h rename to libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl index 36ee849..90bdc2e 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_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(); -- 2.7.4