Add float2half / half2float for F16 inputs/outputs

pwilkin · pwilkin · commit db9eb295cbdf · 2025-09-23T22:06:19.000Z
diff --git a/ggml/src/ggml-cuda/unary.cu b/ggml/src/ggml-cuda/unary.cu
@@ -407,7 +407,7 @@ static __global__ void xielu_kernel(const T * x, T * dst, const int k, float alp
         return;
     }
 
-    const float xi = (float)x[i];
+    const float xi = x->type == GGML_TYPE_F32 ? (float) x[i] : __half2float(x[i]);
     const float gate_pos = (xi > 0.0f);
 
     const float y_pos = alpha_p * xi * xi + beta * xi;
@@ -417,7 +417,7 @@ static __global__ void xielu_kernel(const T * x, T * dst, const int k, float alp
 
     const float out = gate_pos * y_pos + (1.0f - gate_pos) * y_neg;
 
-    dst[i] = (T)out;
+    dst[i] = (T) (dst->type == GGML_TYPE_F32 ? out : __float2half(out));
 }
 
 template <typename T>

Original file line number	Diff line number	Diff line change
`@@ -407,7 +407,7 @@ static __global__ void xielu_kernel(const T * x, T * dst, const int k, float alp`
`407`	`407`	`return;`
`408`	`408`	`}`
`409`	`409`
`410`		`- const float xi = (float)x[i];`
	`410`	`+ const float xi = x->type == GGML_TYPE_F32 ? (float) x[i] : __half2float(x[i]);`
`411`	`411`	`const float gate_pos = (xi > 0.0f);`
`412`	`412`
`413`	`413`	`const float y_pos = alpha_p * xi * xi + beta * xi;`
`@@ -417,7 +417,7 @@ static __global__ void xielu_kernel(const T * x, T * dst, const int k, float alp`
`417`	`417`
`418`	`418`	`const float out = gate_pos * y_pos + (1.0f - gate_pos) * y_neg;`
`419`	`419`
`420`		`- dst[i] = (T)out;`
	`420`	`+ dst[i] = (T) (dst->type == GGML_TYPE_F32 ? out : __float2half(out));`
`421`	`421`	`}`
`422`	`422`
`423`	`423`	`template <typename T>`