--- /dev/null
+_CLC_DECL void barrier(cl_mem_fence_flags flags);
--- /dev/null
+_CLC_DECL size_t get_global_id(uint dim);
--- /dev/null
+_CLC_DECL size_t get_global_size(uint dim);
--- /dev/null
+_CLC_DECL size_t get_group_id(uint dim);
--- /dev/null
+_CLC_DECL size_t get_local_id(uint dim);
--- /dev/null
+_CLC_DECL size_t get_local_size(uint dim);
--- /dev/null
+_CLC_DECL size_t get_num_groups(uint dim);
integer/sub_sat_impl.ll
math/hypot.cl
math/mad.cl
+workitem/get_global_id.cl
+workitem/get_global_size.cl
--- /dev/null
+#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);
+}
--- /dev/null
+#include <clc/clc.h>
+
+_CLC_DEF size_t get_global_size(uint dim) {
+ return get_num_groups(dim)*get_local_size(dim);
+}
+++ /dev/null
-_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;
- }
-}
+++ /dev/null
-_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;
- }
-}
+workitem/get_group_id.cl
+workitem/get_local_id.cl
+workitem/get_local_size.cl
+workitem/get_num_groups.cl
-_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);
}
-_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();
-_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();
-_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();
-_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();