diff --git a/libclc/generic/include/clc/synchronization/barrier.h b/libclc/generic/include/clc/synchronization/barrier.h
new file mode 100644
index 000000000000..7167a3d3f093
--- /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 000000000000..92759f146894
--- /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 000000000000..2f8370585397
--- /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 000000000000..346c82c6c316
--- /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 000000000000..169aeed86786
--- /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 000000000000..040ec58a3d8b
--- /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 000000000000..e555c7efc2d2
--- /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 344c8650fc1f..1d56c4054058 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 000000000000..fdd83d2953d4
--- /dev/null
+++ b/libclc/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/libclc/generic/lib/workitem/get_global_size.cl b/libclc/generic/lib/workitem/get_global_size.cl
new file mode 100644
index 000000000000..5ae649e10d51
--- /dev/null
+++ b/libclc/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/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 026d2fea31f8..000000000000
--- 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 5cd4222d5ae7..000000000000
--- 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 e69de29bb2d1..1a96a1ab8733 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 cd9f3276af62..fb36c2612be4 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
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 18b1bd4db079..2b35b4eaaa95 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
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 1b8c776badf6..f0cfdc005fe8 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
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 cbc1f6ee2893..c3f542595def 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
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 36ee849c613f..90bdc2e41d2c 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();