diff --git a/src/operators/kernel/cl/cl_kernel/softmax.cl b/src/operators/kernel/cl/cl_kernel/softmax.cl index a4514c70640dd6f9582a7362d489f42a58556dcb..215ec69fc283dcb2b538300cb5591b2b9e4b6a13 100644 --- a/src/operators/kernel/cl/cl_kernel/softmax.cl +++ b/src/operators/kernel/cl/cl_kernel/softmax.cl @@ -33,17 +33,17 @@ __kernel void softmax(__read_only image2d_t input_image, maxv = max(maxv, max(temp.x, max(temp.y, max(temp.z, temp.w)))); } - half4 rsum = (half4)(0.0f); + half4 rsum = (half4)(0.0f); for (int i = 0; i < group; ++i) { half4 r = read_imageh(input_image, sampler, (int2)(i, 0)); - rsum += exp(r - maxv); + rsum += convert_half4(exp(convert_float4(r - maxv))); } float sum = rsum.x + rsum.y + rsum.z + rsum.w; half4 rr = read_imageh(input_image, sampler, (int2)(out_w, out_nh)); - half4 result = exp(rr - maxv) / sum; + half4 result = convert_half4(exp(convert_float4(rr - maxv)) / sum); write_imageh(output_image, (int2)(out_w, out_nh), result); }