disable cl_khr_fp16 when __CLPY_ENABLE_CL_KHR_FP16 is not defined

clpy/core/include/cl_stub.hpp

@@ -11,8 +11,10 @@ typedef unsigned long ulong;
 typedef unsigned long size_t;
 typedef long ptrdiff_t;
 
+#ifdef __CLPY_ENABLE_CL_KHR_FP16
 typedef float half;
 typedef half __clpy__half;
+#endif
 #define half __clpy__half
 
 __attribute__((annotate("clpy_no_mangle"))) static unsigned int atomic_cmpxchg(volatile __global unsigned int*, unsigned int, unsigned int);

clpy/core/include/clpy/carray.clh

@@ -1,5 +1,7 @@
 #pragma once
+#ifdef __CLPY_ENABLE_CL_KHR_FP16
 #pragma OPENCL EXTENSION cl_khr_fp16: enable
+#endif
 
 // TODO: Implement common functions in OpenCL C
 #if 0
@@ -530,54 +532,80 @@ static void __clpy_end_print_out() __attribute__((annotate("clpy_end_print_out")
 
 
 #ifdef __ULTIMA
-__attribute__((annotate("clpy_no_mangle"))) static half convert_float_to_half(float x);
+__attribute__((annotate("clpy_no_mangle"))) static ushort convert_float_to_half_ushort(float x);
 #else
 #include "fp16.clh"
+#endif
+
+#ifdef __CLPY_ENABLE_CL_KHR_FP16
+
+#ifndef __ULTIMA
+
 typedef half __clpy__half;
+
 #endif
 
-#ifdef __ULTIMA
-__attribute__((annotate("clpy_no_mangle"))) static half clpy_nextafter_fp16(half x1, half x2);
+__attribute__((annotate("clpy_no_mangle"))) static __clpy__half clpy_nextafter_fp16(__clpy__half x1, __clpy__half x2){
+  return nextafter(x1, x2);
+}
+
+__attribute__((annotate("clpy_no_mangle"))) static __clpy__half convert_float_to_half(float x){
+  const ushort h = convert_float_to_half_ushort(x);
+  return *(const half*)&h;
+}
+
 #else
-static int isnan_fp16(half x){
-  unsigned short const* x_raw = (unsigned short const*)&x;
-  return (*x_raw & 0x7c00u) == 0x7c00u && (*x_raw & 0x03ffu) != 0x0000u;
-}
-static int isfinite_fp16(half x){
-  unsigned short const* x_raw = (unsigned short const*)&x;
-  return (*x_raw & 0x7c00u) != 0x7c00u;
-}
-static int iszero_fp16(half x){
-  unsigned short const* x_raw = (unsigned short const*)&x;
-  return (*x_raw & 0x7fffu) == 0;
-}
-static int eq_nonan_fp16(half x1, half x2){
-  unsigned short const* x1_raw = (unsigned short const*)&x1;
-  unsigned short const* x2_raw = (unsigned short const*)&x2;
-  return (*x1_raw == *x2_raw || ((*x1_raw | *x2_raw) & 0x7fff) == 0);
-}
-static half clpy_nextafter_fp16(half x1, half x2){
-  unsigned short const* x1_raw = (unsigned short const*)&x1;
-  unsigned short const* x2_raw = (unsigned short const*)&x2;
+
+#ifdef __ULTIMA
+static void __clpy_begin_print_out() __attribute__((annotate("clpy_begin_print_out")));
+
+typedef struct __attribute__((packed)) __attribute__((aligned(2))) __clpy__half_{
+  ushort v;
+}__clpy__half;
+
+static int isnan_fp16(__clpy__half x){
+  return (x.v & 0x7c00u) == 0x7c00u && (x.v & 0x03ffu) != 0x0000u;
+}
+static int isfinite_fp16(__clpy__half x){
+  return (x.v & 0x7c00u) != 0x7c00u;
+}
+static int iszero_fp16(__clpy__half x){
+  return (x.v & 0x7fffu) == 0;
+}
+static int eq_nonan_fp16(__clpy__half x1, __clpy__half x2){
+  return (x1.v == x2.v || ((x1.v | x2.v) & 0x7fff) == 0);
+}
+__attribute__((annotate("clpy_no_mangle"))) static __clpy__half clpy_nextafter_fp16(__clpy__half x1, __clpy__half x2){
   unsigned short ret_raw_;
 
   if (!isfinite_fp16(x1) || isnan_fp16(x2)){
     ret_raw_ = 0x7e00u; // NaN in fp16
   }else if(eq_nonan_fp16(x1, x2)){
-    ret_raw_ = *x1_raw;
+    ret_raw_ = x1.v;
   }else if(iszero_fp16(x1)){
-    ret_raw_ = (*x2_raw & 0x8000u) + 1;
-  }else if(!(*x1_raw & 0x8000u)){
-    if (*(short const*)x1_raw > *(short const*)x2_raw){
-      ret_raw_ = *x1_raw - 1;
+    ret_raw_ = (x2.v & 0x8000u) + 1;
+  }else if(!(x1.v & 0x8000u)){
+    if (*(short const*)&x1.v > *(short const*)&x2.v){
+      ret_raw_ = x1.v - 1;
     }else{
-      ret_raw_ = *x1_raw + 1;
+      ret_raw_ = x1.v + 1;
     }
-  }else if(!(*x2_raw & 0x8000u) || (*x1_raw & 0x7fffu) > (*x2_raw & 0x7fffu)) {
-    ret_raw_ = *x1_raw - 1;
+  }else if(!(x2.v & 0x8000u) || (x1.v & 0x7fffu) > (x2.v & 0x7fffu)) {
+    ret_raw_ = x1.v - 1;
   } else {
-    ret_raw_ = *x1_raw + 1;
+    ret_raw_ = x1.v + 1;
   }
-  return *(half*)&ret_raw_;
+  __clpy__half ret;
+  ret.v = ret_raw_;
+  return ret;
+}
+
+__attribute__((annotate("clpy_no_mangle"))) static __clpy__half convert_float_to_half(float x){
+  __clpy__half ret;
+  ret.v = convert_float_to_half_ushort(x);
+  return ret;
 }
+
+static void __clpy_end_print_out() __attribute__((annotate("clpy_end_print_out")));
+#endif
 #endif

clpy/core/include/clpy/fp16.clh

@@ -12,7 +12,7 @@
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
  */
 
-static half convert_float_to_half(float f) {
+static ushort convert_float_to_half_ushort(float f) {
   const float scale_to_inf = 0x1.0p+112f;
   const float scale_to_zero = 0x1.0p-110f;
   float base = (fabs(f) * scale_to_inf) * scale_to_zero;
@@ -32,5 +32,5 @@ static half convert_float_to_half(float f) {
   const uint mantissa_bits = bits & (uint)(0x00000FFF);
   const uint nonsign = exp_bits + mantissa_bits;
   const ushort ret = ((sign >> 16) | (shl1_w > (uint)(0xFF000000) ? (ushort)(0x7E00) : nonsign));
-  return *(const half*)&ret;
+  return ret;
 }