[libclc] 21/291: PTX: move implementations of work-item and synchronisation functions to lib, and add header files in generic. Incorporates a patch by Tom Stellard!
Andreas Beckmann
anbe at moszumanska.debian.org
Tue Sep 8 10:53:28 UTC 2015
This is an automated email from the git hooks/post-receive script.
anbe pushed a commit to branch master
in repository libclc.
commit a37e2c06b60184ad139cf01f6b37ef2276febf34
Author: Peter Collingbourne <peter at pcc.me.uk>
Date: Sun Aug 5 22:25:37 2012 +0000
PTX: move implementations of work-item and synchronisation functions
to lib, and add header files in generic. Incorporates a patch by
Tom Stellard!
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@161313 91177308-0d34-0410-b5e6-96231b3b80d8
---
generic/include/clc/synchronization/barrier.h | 1 +
generic/include/clc/workitem/get_global_id.h | 1 +
generic/include/clc/workitem/get_global_size.h | 1 +
generic/include/clc/workitem/get_group_id.h | 1 +
generic/include/clc/workitem/get_local_id.h | 1 +
generic/include/clc/workitem/get_local_size.h | 1 +
generic/include/clc/workitem/get_num_groups.h | 1 +
generic/lib/SOURCES | 2 ++
generic/lib/workitem/get_global_id.cl | 5 +++++
generic/lib/workitem/get_global_size.cl | 5 +++++
ptx-nvidiacl/include/clc/workitem/get_global_id.h | 8 --------
ptx-nvidiacl/include/clc/workitem/get_global_size.h | 8 --------
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(-)
diff --git a/generic/include/clc/synchronization/barrier.h b/generic/include/clc/synchronization/barrier.h
new file mode 100644
index 0000000..7167a3d
--- /dev/null
+++ b/generic/include/clc/synchronization/barrier.h
@@ -0,0 +1 @@
+_CLC_DECL void barrier(cl_mem_fence_flags flags);
diff --git a/generic/include/clc/workitem/get_global_id.h b/generic/include/clc/workitem/get_global_id.h
new file mode 100644
index 0000000..92759f1
--- /dev/null
+++ b/generic/include/clc/workitem/get_global_id.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_id(uint dim);
diff --git a/generic/include/clc/workitem/get_global_size.h b/generic/include/clc/workitem/get_global_size.h
new file mode 100644
index 0000000..2f83705
--- /dev/null
+++ b/generic/include/clc/workitem/get_global_size.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_size(uint dim);
diff --git a/generic/include/clc/workitem/get_group_id.h b/generic/include/clc/workitem/get_group_id.h
new file mode 100644
index 0000000..346c82c
--- /dev/null
+++ b/generic/include/clc/workitem/get_group_id.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_group_id(uint dim);
diff --git a/generic/include/clc/workitem/get_local_id.h b/generic/include/clc/workitem/get_local_id.h
new file mode 100644
index 0000000..169aeed
--- /dev/null
+++ b/generic/include/clc/workitem/get_local_id.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_local_id(uint dim);
diff --git a/generic/include/clc/workitem/get_local_size.h b/generic/include/clc/workitem/get_local_size.h
new file mode 100644
index 0000000..040ec58
--- /dev/null
+++ b/generic/include/clc/workitem/get_local_size.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_local_size(uint dim);
diff --git a/generic/include/clc/workitem/get_num_groups.h b/generic/include/clc/workitem/get_num_groups.h
new file mode 100644
index 0000000..e555c7e
--- /dev/null
+++ b/generic/include/clc/workitem/get_num_groups.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_num_groups(uint dim);
diff --git a/generic/lib/SOURCES b/generic/lib/SOURCES
index 344c865..1d56c40 100644
--- a/generic/lib/SOURCES
+++ b/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/generic/lib/workitem/get_global_id.cl b/generic/lib/workitem/get_global_id.cl
new file mode 100644
index 0000000..fdd83d2
--- /dev/null
+++ b/generic/lib/workitem/get_global_id.cl
@@ -0,0 +1,5 @@
+#include <clc/clc.h>
+
+_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/generic/lib/workitem/get_global_size.cl b/generic/lib/workitem/get_global_size.cl
new file mode 100644
index 0000000..5ae649e
--- /dev/null
+++ b/generic/lib/workitem/get_global_size.cl
@@ -0,0 +1,5 @@
+#include <clc/clc.h>
+
+_CLC_DEF size_t get_global_size(uint dim) {
+ return get_num_groups(dim)*get_local_size(dim);
+}
diff --git a/ptx-nvidiacl/include/clc/workitem/get_global_id.h b/ptx-nvidiacl/include/clc/workitem/get_global_id.h
deleted file mode 100644
index 026d2fe..0000000
--- a/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/ptx-nvidiacl/include/clc/workitem/get_global_size.h b/ptx-nvidiacl/include/clc/workitem/get_global_size.h
deleted file mode 100644
index 5cd4222..0000000
--- a/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/ptx-nvidiacl/lib/SOURCES b/ptx-nvidiacl/lib/SOURCES
index e69de29..1a96a1a 100644
--- a/ptx-nvidiacl/lib/SOURCES
+++ b/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/ptx-nvidiacl/include/clc/synchronization/barrier.h b/ptx-nvidiacl/lib/synchronization/barrier.cl
similarity index 51%
rename from ptx-nvidiacl/include/clc/synchronization/barrier.h
rename to ptx-nvidiacl/lib/synchronization/barrier.cl
index cd9f327..fb36c26 100644
--- a/ptx-nvidiacl/include/clc/synchronization/barrier.h
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_group_id.h b/ptx-nvidiacl/lib/workitem/get_group_id.cl
similarity index 74%
rename from ptx-nvidiacl/include/clc/workitem/get_group_id.h
rename to ptx-nvidiacl/lib/workitem/get_group_id.cl
index 18b1bd4..2b35b4e 100644
--- a/ptx-nvidiacl/include/clc/workitem/get_group_id.h
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_local_id.h b/ptx-nvidiacl/lib/workitem/get_local_id.cl
similarity index 74%
rename from ptx-nvidiacl/include/clc/workitem/get_local_id.h
rename to ptx-nvidiacl/lib/workitem/get_local_id.cl
index 1b8c776..f0cfdc0 100644
--- a/ptx-nvidiacl/include/clc/workitem/get_local_id.h
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_local_size.h b/ptx-nvidiacl/lib/workitem/get_local_size.cl
similarity index 74%
rename from ptx-nvidiacl/include/clc/workitem/get_local_size.h
rename to ptx-nvidiacl/lib/workitem/get_local_size.cl
index cbc1f6e..c3f5425 100644
--- a/ptx-nvidiacl/include/clc/workitem/get_local_size.h
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_num_groups.h b/ptx-nvidiacl/lib/workitem/get_num_groups.cl
similarity index 74%
rename from ptx-nvidiacl/include/clc/workitem/get_num_groups.h
rename to ptx-nvidiacl/lib/workitem/get_num_groups.cl
index 36ee849..90bdc2e 100644
--- a/ptx-nvidiacl/include/clc/workitem/get_num_groups.h
+++ b/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();
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git
More information about the Pkg-opencl-commits
mailing list