Initial commit.

git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@147756 91177308-0d34-0410-b5e6-96231b3b80d8

8 files changed, 54 insertions, 0 deletions

diff --git a/ptx-nvidiacl/include/clc/synchronization/barrier.h b/ptx-nvidiacl/include/clc/synchronization/barrier.h
new file mode 100644
index 0000000..cd9f327
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_global_id.h b/ptx-nvidiacl/include/clc/workitem/get_global_id.h
new file mode 100644
index 0000000..026d2fe
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_global_size.h b/ptx-nvidiacl/include/clc/workitem/get_global_size.h
new file mode 100644
index 0000000..5cd4222
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_group_id.h b/ptx-nvidiacl/include/clc/workitem/get_group_id.h
new file mode 100644
index 0000000..18b1bd4
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_local_id.h b/ptx-nvidiacl/include/clc/workitem/get_local_id.h
new file mode 100644
index 0000000..1b8c776
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_local_size.h b/ptx-nvidiacl/include/clc/workitem/get_local_size.h
new file mode 100644
index 0000000..cbc1f6e
--- /dev/null
+++ b/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/ptx-nvidiacl/include/clc/workitem/get_num_groups.h b/ptx-nvidiacl/include/clc/workitem/get_num_groups.h
new file mode 100644
index 0000000..36ee849
--- /dev/null
+++ b/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/ptx-nvidiacl/lib/SOURCES b/ptx-nvidiacl/lib/SOURCES
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/ptx-nvidiacl/lib/SOURCES