opencl: do sqrt in f32 then convert to f16 for better precision

lhez · lhez · commit 338d9acb10a4 · 2025-11-24T22:28:49.000-08:00
diff --git a/ggml/src/ggml-opencl/kernels/sqrt.cl b/ggml/src/ggml-opencl/kernels/sqrt.cl
@@ -36,7 +36,7 @@ kernel void kernel_sqrt_cont_f16(
     dst  = (global half*)((global char*)dst + offsetd);
 
     uint gid = get_global_id(0);
-    dst[gid] = sqrt(src0[gid]);
+    dst[gid] = convert_half(sqrt(convert_float(src0[gid])));
 }
 
 kernel void kernel_sqrt_cont_f16_4(
@@ -49,5 +49,5 @@ kernel void kernel_sqrt_cont_f16_4(
     dst  = (global half4*)((global char*)dst + offsetd);
 
     uint gid = get_global_id(0);
-    dst[gid] = sqrt(src0[gid]);
+    dst[gid] = convert_half4(sqrt(convert_float4(src0[gid])));
 }

Original file line number	Diff line number	Diff line change
`@@ -36,7 +36,7 @@ kernel void kernel_sqrt_cont_f16(`
`36`	`36`	`dst = (global half)((global char)dst + offsetd);`
`37`	`37`
`38`	`38`	`uint gid = get_global_id(0);`
`39`		`- dst[gid] = sqrt(src0[gid]);`
	`39`	`+ dst[gid] = convert_half(sqrt(convert_float(src0[gid])));`
`40`	`40`	`}`
`41`	`41`
`42`	`42`	`kernel void kernel_sqrt_cont_f16_4(`
`@@ -49,5 +49,5 @@ kernel void kernel_sqrt_cont_f16_4(`
`49`	`49`	`dst = (global half4)((global char)dst + offsetd);`
`50`	`50`
`51`	`51`	`uint gid = get_global_id(0);`
`52`		`- dst[gid] = sqrt(src0[gid]);`
	`52`	`+ dst[gid] = convert_half4(sqrt(convert_float4(src0[gid])));`
`53`	`53`	`}`