diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-08-08 11:31:27 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-08-12 18:16:56 +0800 |
| commit | 4889cdf76678eeda085e2bfc49c7a57e2ee0a6f9 (patch) | |
| tree | 2c7b0bb802e0d8f5fcae2cca1acf26bd4fbb6e86 /kernels/compiler_math_3op.cl | |
| parent | b673343987f751b8c006cc4920111545073827cb (diff) | |
Utest: Add half type mad test case
Mad now can support half type, add a test.
Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Diffstat (limited to 'kernels/compiler_math_3op.cl')
| -rw-r--r-- | kernels/compiler_math_3op.cl | 17 |
1 files changed, 16 insertions, 1 deletions
diff --git a/kernels/compiler_math_3op.cl b/kernels/compiler_math_3op.cl index 9194ef01..1a43e1bc 100644 --- a/kernels/compiler_math_3op.cl +++ b/kernels/compiler_math_3op.cl @@ -1,4 +1,18 @@ -kernel void compiler_math_3op(global float *dst, global float *src1, global float *src2, global float *src3) { +#ifdef HALF +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +kernel void compiler_math_3op_half(global half *dst, global half *src1, global half *src2, global half *src3) { + int i = get_global_id(0); + const half x = src1[i], y = src2[i], z = src3[i]; + switch (i%2) { + case 0: dst[i] = mad(x, y, z); break; + case 1: dst[i] = fma(x, y, z); break; + default: dst[i] = 1.f; break; + }; + dst[0] = mad(src1[0],src2[0],src3[0]); +} +#else +kernel void compiler_math_3op_float(global float *dst, global float *src1, global float *src2, global float *src3) { + int i = get_global_id(0); const float x = src1[i], y = src2[i], z = src3[i]; switch (i%2) { @@ -8,3 +22,4 @@ kernel void compiler_math_3op(global float *dst, global float *src1, global floa }; dst[0] = mad(src1[0],src2[0],src3[0]); } +#endif |