From 0cf109f6566b7b9bab915549a81a22b370011c34 Mon Sep 17 00:00:00 2001
From: doe300 <stadeldani@web.de>
Date: Sun, 25 Feb 2018 22:15:45 +0100
Subject: [PATCH] Fixes missing float conversion functions

---
 include/_async.h       |  1 -
 include/_conversions.h | 12 ++++++------
 include/_vector.h      |  3 +++
 include/_work_items.h  | 10 ++++++++++
 4 files changed, 19 insertions(+), 7 deletions(-)

diff --git a/include/_async.h b/include/_async.h
index 4b4acba..2a88518 100644
--- a/include/_async.h
+++ b/include/_async.h
@@ -159,7 +159,6 @@
 			return event; \
 		}
 
-//Since it doesn't affect the functional behavior, the implementation is a no-op
 #define PREFETCH(type) \
 		INLINE void prefetch(const __global type * ptr, size_t num_entries) OVERLOADABLE \
 		{ \
diff --git a/include/_conversions.h b/include/_conversions.h
index 021c8a9..0560ddd 100644
--- a/include/_conversions.h
+++ b/include/_conversions.h
@@ -120,27 +120,27 @@
 		}
 
 #define CONVERT_FLOAT_TO_FLOAT(saturation, rounding) \
-		INLINE float convert_##destType##saturation##rounding(float val) OVERLOADABLE CONST \
+		INLINE float convert_float##saturation##rounding(float val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		} \
-		INLINE float##2 convert_##destType##2##saturation##rounding(float##2 val) OVERLOADABLE CONST \
+		INLINE float##2 convert_float2##saturation##rounding(float##2 val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		} \
-		INLINE float##3 convert_##destType##3##saturation##rounding(float##3 val) OVERLOADABLE CONST \
+		INLINE float##3 convert_float3##saturation##rounding(float##3 val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		} \
-		INLINE float##4 convert_##destType##4##saturation##rounding(float##4 val) OVERLOADABLE CONST \
+		INLINE float##4 convert_float4##saturation##rounding(float##4 val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		} \
-		INLINE float##8 convert_##destType##8##saturation##rounding(float##8 val) OVERLOADABLE CONST \
+		INLINE float##8 convert_float8##saturation##rounding(float##8 val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		} \
-		INLINE float##16 convert_##destType##16##saturation##rounding(float##16 val) OVERLOADABLE CONST \
+		INLINE float##16 convert_float16##saturation##rounding(float##16 val) OVERLOADABLE CONST \
 		{ \
 			return ROUND_TO_INTEGER(rounding, val); \
 		}
diff --git a/include/_vector.h b/include/_vector.h
index a17c691..4f90f95 100644
--- a/include/_vector.h
+++ b/include/_vector.h
@@ -218,6 +218,9 @@ VECTOR_STORE(uint)
 VECTOR_STORE(int)
 VECTOR_STORE(float)
 
+//TODO vload(a)_half, vload(a)_halfn (+rounding) (load half and return converted to float, possible with unpack-modes)
+//TODO vstore(a)_half, vstore(a)_halfn (+rounding) (store float as half in memory, possible with pack modes)
+
 /*
  * TODO shuffle2, but LLVM fails, since the indices for the __builtin intrinsic need to be constant integers!
 VECTOR_SHUFFLE_2(uchar, uchar)
diff --git a/include/_work_items.h b/include/_work_items.h
index 8f949d9..53e57ba 100644
--- a/include/_work_items.h
+++ b/include/_work_items.h
@@ -63,6 +63,16 @@ INLINE size_t get_local_linear_id() OVERLOADABLE CONST
 	return vc4cl_mul24(vc4cl_mul24(vc4cl_local_id(2), vc4cl_local_size(1), VC4CL_UNSIGNED), vc4cl_bitcast_uint(vc4cl_local_size(0)), VC4CL_UNSIGNED) + vc4cl_mul24(vc4cl_local_id(1), vc4cl_local_size(0), VC4CL_UNSIGNED) + vc4cl_local_id(0);
 }
 
+/*
+ * Returns the work-items 1-dimensional global ID.
+ */
+size_t get_global_linear_id(void) CONST; // prototype
+size_t get_global_linear_id(void) OVERLOADABLE CONST
+{
+	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));
+}
 
 #endif /* VC4CL_WORK_ITEMS_H */