Fixes missing float conversion functions

doe300 · Feb 25, 2018 · 0cf109f · 0cf109f
1 parent 8ba7862
commit 0cf109f
Show file tree

Hide file tree

Showing 4 changed files with 19 additions and 7 deletions.
diff --git 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
@@ -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
@@ -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
@@ -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 */