[Beignet] [PATCH V2 2/2] enable utest compiler_math_3op for mad test.

Guo, Yejun yejun.guo at intel.com
Wed Apr 27 07:52:37 UTC 2016


this patch set looks good to me, thanks

-----Original Message-----
From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of xionghu.luo at intel.com
Sent: Wednesday, April 27, 2016 11:29 PM
To: beignet at lists.freedesktop.org
Cc: Luo, Xionghu
Subject: [Beignet] [PATCH V2 2/2] enable utest compiler_math_3op for mad test.

From: Luo Xionghu <xionghu.luo at intel.com>

 v2: add uniform dest test coverage.
Signed-off-by: Luo Xionghu <xionghu.luo at intel.com>
---
 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
Beignet at lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list