[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