Fix instrinsic type casting issues (#145)

microsoft · Oct 13, 2023 · bde1df9 · bde1df9
1 parent a0e2d2c
commit bde1df9
Showing 1 changed file with 2 additions and 2 deletions.
diff --git a/ark/include/kernels/type_intrinsics.h b/ark/include/kernels/type_intrinsics.h
@@ -91,7 +91,7 @@ struct Div {
 
 struct Exp {
     static DEVICE float compute(float input) { return expf(input); }
-    static DEVICE __half compute(__half input) { return hexp(input); }
+    static DEVICE half compute(half input) { return half(expf(float(input))); }
     static DEVICE bfloat16 compute(bfloat16 input) {
         return bfloat16(expf(float(input)));
     }
@@ -103,7 +103,7 @@ struct Exp {
 
 struct Sqrt {
     static DEVICE float compute(float input) { return sqrtf(input); }
-    static DEVICE __half compute(__half input) { return hsqrt(input); }
+    static DEVICE half compute(half input) { return half(sqrtf(float(input))); }
     static DEVICE bfloat16 compute(bfloat16 input) {
         return bfloat16(sqrtf(float(input)));
     }