this patch set looks good to me, thanks -----Original Message----- From: Beignet [mailto:[email protected]] On Behalf Of [email protected] Sent: Wednesday, April 27, 2016 11:29 PM To: [email protected] Cc: Luo, Xionghu Subject: [Beignet] [PATCH V2 2/2] enable utest compiler_math_3op for mad test.
From: Luo Xionghu <[email protected]> v2: add uniform dest test coverage. Signed-off-by: Luo Xionghu <[email protected]> --- kernels/compiler_math_3op.cl | 3 ++- utests/CMakeLists.txt | 1 + utests/compiler_math_3op.cpp | 10 ++++++---- 3 files changed, 9 insertions(+), 5 deletions(-) diff --git a/kernels/compiler_math_3op.cl b/kernels/compiler_math_3op.cl index 95b0398..9194ef0 100644 --- a/kernels/compiler_math_3op.cl +++ b/kernels/compiler_math_3op.cl @@ -1,9 +1,10 @@ kernel void compiler_math_3op(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) { + 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]); } diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 4742c4a..a76bc47 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -268,6 +268,7 @@ set (utests_sources builtin_global_linear_id.cpp builtin_local_linear_id.cpp compiler_mix.cpp + compiler_math_3op.cpp compiler_bsort.cpp) if (LLVM_VERSION_NODOT VERSION_GREATER 34) diff --git a/utests/compiler_math_3op.cpp b/utests/compiler_math_3op.cpp index f90f9d6..523b72b 100644 --- a/utests/compiler_math_3op.cpp +++ b/utests/compiler_math_3op.cpp @@ -5,11 +5,12 @@ static void cpu_compiler_math(float *dst, float *src1, float *src2, float *src3, int i) { const float x = src1[i], y = src2[i], z = src3[i]; - switch (i) { + switch (i%2) { case 0: dst[i] = x * y + z; break; case 1: dst[i] = x * y + z; break; default: dst[i] = 1.f; break; }; + dst[0] = (src1[0]*src2[0]+src3[0]); } static void compiler_math_3op(void) @@ -35,9 +36,9 @@ static void compiler_math_3op(void) OCL_MAP_BUFFER(2); OCL_MAP_BUFFER(3); for (uint32_t i = 0; i < 32; ++i) { - cpu_src1[i] = ((float*)buf_data[1])[i] = .1f * (rand() & 15); - cpu_src2[i] = ((float*)buf_data[2])[i] = .1f * (rand() & 15); - cpu_src3[i] = ((float*)buf_data[3])[i] = .1f * (rand() & 15); + cpu_src1[i] = ((float*)buf_data[1])[i] = .001f * (rand() & 15); + cpu_src2[i] = ((float*)buf_data[2])[i] = .002f * (rand() & 15); + cpu_src3[i] = ((float*)buf_data[3])[i] = .003f * (rand() & 15); } OCL_UNMAP_BUFFER(1); OCL_UNMAP_BUFFER(2); @@ -50,6 +51,7 @@ static void compiler_math_3op(void) for (int i = 0; i < 16; ++i) { const float cpu = cpu_dst[i]; const float gpu = ((float*)buf_data[0])[i]; + //printf("cpu:%f, gpu:%f\n", cpu, gpu); if (std::isinf(cpu)) OCL_ASSERT(std::isinf(gpu)); else if (std::isnan(cpu)) -- 2.1.4 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
