[cfe-commits] [libclc] r161313 - in /libclc/trunk: generic/include/clc/synchronization/ generic/include/clc/workitem/ generic/lib/ generic/lib/workitem/ ptx-nvidiacl/include/clc/synchronization/ ptx-nvidiacl/include/clc/workitem/ ptx-nvidiacl/lib/ ptx-nvidiacl/lib/synchronization/ ptx-nvidiacl/lib/workitem/
Peter Collingbourne
peter at pcc.me.uk
Sun Aug 5 15:25:38 PDT 2012
Author: pcc
Date: Sun Aug 5 17:25:37 2012
New Revision: 161313
URL: http://llvm.org/viewvc/llvm-project?rev=161313&view=rev
Log:
PTX: move implementations of work-item and synchronisation functions
to lib, and add header files in generic. Incorporates a patch by
Tom Stellard!
Added:
libclc/trunk/generic/include/clc/synchronization/barrier.h
libclc/trunk/generic/include/clc/workitem/
libclc/trunk/generic/include/clc/workitem/get_global_id.h
libclc/trunk/generic/include/clc/workitem/get_global_size.h
libclc/trunk/generic/include/clc/workitem/get_group_id.h
libclc/trunk/generic/include/clc/workitem/get_local_id.h
libclc/trunk/generic/include/clc/workitem/get_local_size.h
libclc/trunk/generic/include/clc/workitem/get_num_groups.h
libclc/trunk/generic/lib/workitem/
libclc/trunk/generic/lib/workitem/get_global_id.cl
libclc/trunk/generic/lib/workitem/get_global_size.cl
libclc/trunk/ptx-nvidiacl/lib/synchronization/
libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl
- copied, changed from r161312, libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h
libclc/trunk/ptx-nvidiacl/lib/workitem/
libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl
- copied, changed from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h
libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl
- copied, changed from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h
libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl
- copied, changed from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h
libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl
- copied, changed from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h
Removed:
libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_id.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_size.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h
libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h
Modified:
libclc/trunk/generic/lib/SOURCES
libclc/trunk/ptx-nvidiacl/lib/SOURCES
Added: libclc/trunk/generic/include/clc/synchronization/barrier.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/synchronization/barrier.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/synchronization/barrier.h (added)
+++ libclc/trunk/generic/include/clc/synchronization/barrier.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL void barrier(cl_mem_fence_flags flags);
Added: libclc/trunk/generic/include/clc/workitem/get_global_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_global_id.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_global_id.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_global_id.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_id(uint dim);
Added: libclc/trunk/generic/include/clc/workitem/get_global_size.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_global_size.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_global_size.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_global_size.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_size(uint dim);
Added: libclc/trunk/generic/include/clc/workitem/get_group_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_group_id.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_group_id.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_group_id.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_group_id(uint dim);
Added: libclc/trunk/generic/include/clc/workitem/get_local_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_local_id.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_local_id.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_local_id.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_local_id(uint dim);
Added: libclc/trunk/generic/include/clc/workitem/get_local_size.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_local_size.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_local_size.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_local_size.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_local_size(uint dim);
Added: libclc/trunk/generic/include/clc/workitem/get_num_groups.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_num_groups.h?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_num_groups.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_num_groups.h Sun Aug 5 17:25:37 2012
@@ -0,0 +1 @@
+_CLC_DECL size_t get_num_groups(uint dim);
Modified: libclc/trunk/generic/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/SOURCES?rev=161313&r1=161312&r2=161313&view=diff
==============================================================================
--- libclc/trunk/generic/lib/SOURCES (original)
+++ libclc/trunk/generic/lib/SOURCES Sun Aug 5 17:25:37 2012
@@ -12,3 +12,5 @@
integer/sub_sat_impl.ll
math/hypot.cl
math/mad.cl
+workitem/get_global_id.cl
+workitem/get_global_size.cl
Added: libclc/trunk/generic/lib/workitem/get_global_id.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/workitem/get_global_id.cl?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/lib/workitem/get_global_id.cl (added)
+++ libclc/trunk/generic/lib/workitem/get_global_id.cl Sun Aug 5 17:25:37 2012
@@ -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);
+}
Added: libclc/trunk/generic/lib/workitem/get_global_size.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/workitem/get_global_size.cl?rev=161313&view=auto
==============================================================================
--- libclc/trunk/generic/lib/workitem/get_global_size.cl (added)
+++ libclc/trunk/generic/lib/workitem/get_global_size.cl Sun Aug 5 17:25:37 2012
@@ -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);
+}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h (removed)
@@ -1,6 +0,0 @@
-_CLC_INLINE void barrier(cl_mem_fence_flags flags) {
- if (flags & CLK_LOCAL_MEM_FENCE) {
- __builtin_ptx_bar_sync(0);
- }
-}
-
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_id.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_id.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_id.h (removed)
@@ -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;
- }
-}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_size.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_size.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_size.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_global_size.h (removed)
@@ -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;
- }
-}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h (removed)
@@ -1,8 +0,0 @@
-_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;
- }
-}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h (removed)
@@ -1,8 +0,0 @@
-_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;
- }
-}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h (removed)
@@ -1,8 +0,0 @@
-_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;
- }
-}
Removed: libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h?rev=161312&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h (original)
+++ libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h (removed)
@@ -1,8 +0,0 @@
-_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;
- }
-}
Modified: libclc/trunk/ptx-nvidiacl/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/SOURCES?rev=161313&r1=161312&r2=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/SOURCES (original)
+++ libclc/trunk/ptx-nvidiacl/lib/SOURCES Sun Aug 5 17:25:37 2012
@@ -0,0 +1,4 @@
+workitem/get_group_id.cl
+workitem/get_local_id.cl
+workitem/get_local_size.cl
+workitem/get_num_groups.cl
Copied: libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl (from r161312, libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl?p2=libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl&p1=libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h&r1=161312&r2=161313&rev=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/synchronization/barrier.h (original)
+++ libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl Sun Aug 5 17:25:37 2012
@@ -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);
}
Copied: libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl (from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl?p2=libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl&p1=libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h&r1=161312&r2=161313&rev=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_group_id.h (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl Sun Aug 5 17:25:37 2012
@@ -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();
Copied: libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl (from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl?p2=libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl&p1=libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h&r1=161312&r2=161313&rev=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_id.h (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl Sun Aug 5 17:25:37 2012
@@ -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();
Copied: libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl (from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl?p2=libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl&p1=libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h&r1=161312&r2=161313&rev=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_local_size.h (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl Sun Aug 5 17:25:37 2012
@@ -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();
Copied: libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl (from r161312, libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl?p2=libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl&p1=libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h&r1=161312&r2=161313&rev=161313&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/include/clc/workitem/get_num_groups.h (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl Sun Aug 5 17:25:37 2012
@@ -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();
More information about the cfe-commits
mailing list