Add get global/local linear id by calculate with global/local
id, size and offset. The get_queue_local_size() and get_loal_size()
should be different when the global work group size is not uniform,
but now they are the same. We will refine these functions when we
support non-uniform work-group size.
Signed-off-by: Pan Xiuli <xiuli....@intel.com>
---
backend/src/libocl/include/ocl_workitem.h | 3 +++
backend/src/libocl/src/ocl_workitem.cl | 30 ++++++++++++++++++++++++++++++
2 files changed, 33 insertions(+)
diff --git a/backend/src/libocl/include/ocl_workitem.h
b/backend/src/libocl/include/ocl_workitem.h
index 84bb1fb..c3b0bdb 100644
--- a/backend/src/libocl/include/ocl_workitem.h
+++ b/backend/src/libocl/include/ocl_workitem.h
@@ -24,9 +24,12 @@ OVERLOADABLE uint get_work_dim(void);
OVERLOADABLE uint get_global_size(uint dimindx);
OVERLOADABLE uint get_global_id(uint dimindx);
OVERLOADABLE uint get_local_size(uint dimindx);
+OVERLOADABLE uint get_enqueued_local_size(uint dimindx);
OVERLOADABLE uint get_local_id(uint dimindx);
OVERLOADABLE uint get_num_groups(uint dimindx);
OVERLOADABLE uint get_group_id(uint dimindx);
OVERLOADABLE uint get_global_offset(uint dimindx);
+OVERLOADABLE uint get_global_linear_id(void);
+OVERLOADABLE uint get_local_linear_id(void);
#endif /* __OCL_WORKITEM_H__ */
diff --git a/backend/src/libocl/src/ocl_workitem.cl
b/backend/src/libocl/src/ocl_workitem.cl
index 6ddc406..235f12b 100644
--- a/backend/src/libocl/src/ocl_workitem.cl
+++ b/backend/src/libocl/src/ocl_workitem.cl
@@ -55,3 +55,33 @@ DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)
OVERLOADABLE uint get_global_id(uint dim) {
return get_local_id(dim) + get_local_size(dim) * get_group_id(dim) +
get_global_offset(dim);
}
+
+OVERLOADABLE uint get_enqueued_local_size (uint dimindx)
+{
+ //TODO: should be different with get_local_size when support
+ //non-uniform work-group size
+ return get_local_size(dimindx);
+}
+
+OVERLOADABLE uint get_global_linear_id(void)
+{
+ uint dim = __gen_ocl_get_work_dim();
+ if (dim == 1) return get_global_id(0) - get_global_offset(0);
+ else if (dim == 2) return (get_global_id(1) - get_global_offset(1)) *
get_global_size(0) +
+ get_global_id(0) -get_global_offset(0);
+ else if (dim == 3) return ((get_global_id(2) - get_global_offset(2)) *
+ get_global_size(1) * get_global_size(0)) +
+ ((get_global_id(1) - get_global_offset(1)) *
get_global_size (0)) +
+ (get_global_id(0) - get_global_offset(0));
+ else return 0;
+}
+
+OVERLOADABLE uint get_local_linear_id(void)
+{
+ uint dim = __gen_ocl_get_work_dim();
+ if (dim == 1) return get_local_id(0);
+ else if (dim == 2) return get_local_id(1) * get_local_size (0) +
get_local_id(0);
+ else if (dim == 3) return (get_local_id(2) * get_local_size(1) *
get_local_size(0)) +
+ (get_local_id(1) * get_local_size(0)) +
get_local_id(0);
+ else return 0;
+}
--
2.1.4
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
