fix cuda half math function is undefined: hpow, htanh (apache#6253)

trevor-m · Aug 26, 2020 · 70eed47 · 70eed47
1 parent 0c18547
commit 70eed47
Showing 1 changed file with 16 additions and 0 deletions.
diff --git a/src/target/source/literal/cuda_half_t.h b/src/target/source/literal/cuda_half_t.h
@@ -293,6 +293,22 @@ __pack_half2(const half x, const half y) {
   unsigned v1 = *((unsigned short *)&y);
   return (v1 << 16) | v0;
 }
+
+// fix undefined fp16 match function
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
+static inline __device__ __host__ half hpow(half x, half y) {
+  float tmp_x = __half2float(x);
+  float tmp_y = __half2float(y);
+  float result = powf(tmp_x, tmp_y);
+  return __float2half(result);
+}
+
+static inline __device__ __host__ half htanh(half x) {
+  float tmp_x = __half2float(x);
+  float result = tanhf(tmp_x);
+  return __float2half(result);
+}
+#endif
 )";
 
 static constexpr const char* _cuda_warp_intrinsic_util = R"(