[Beignet] [PATCH] support converting 64-bit integer to 32-bit float

[Homer Hsing]

Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
 backend/src/backend/gen_context.cpp        | 14 ++++++++++
 backend/src/backend/gen_insn_selection.cpp |  4 +++
 backend/src/backend/gen_insn_selection.hxx |  1 +
 kernels/compiler_long_convert.cl           |  5 ++++
 utests/compiler_long_convert.cpp           | 41 ++++++++++++++++++++++++++++++
 5 files changed, 65 insertions(+)

diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
index 3e3a8fc..0b61c0a 100644
--- a/backend/src/backend/gen_context.cpp
+++ b/backend/src/backend/gen_context.cpp
@@ -189,6 +189,20 @@ namespace gbe
       case SEL_OP_MOV_DF:
         p->MOV_DF(dst, src, tmp);
         break;
+      case SEL_OP_CONVI64_TO_F:
+       {
+        tmp.type = src.is_signed_int() ? GEN_TYPE_D : GEN_TYPE_UD;
+        GenRegister tmp2 = GenRegister::suboffset(tmp, p->curr.execWidth);
+        tmp2.type = GEN_TYPE_F;
+        loadTopHalf(tmp, src);
+        p->MOV(dst, tmp);
+        p->MUL(dst, dst, GenRegister::immf(65536.f * 65536.f));
+        tmp.type = GEN_TYPE_UD;
+        loadBottomHalf(tmp, src);
+        p->MOV(tmp2, tmp);
+        p->ADD(dst, dst, tmp2);
+        break;
+       }
       case SEL_OP_CONVI_TO_I64: {
         GenRegister middle;
         if (src.type == GEN_TYPE_B || src.type == GEN_TYPE_D) {
diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
index f1b85bb..94d0993 100644
--- a/backend/src/backend/gen_insn_selection.cpp
+++ b/backend/src/backend/gen_insn_selection.cpp
@@ -460,6 +460,7 @@ namespace gbe
     ALU2(UPSAMPLE_LONG)
     ALU1WithTemp(CONVI_TO_I64)
     ALU1(CONVI64_TO_I)
+    ALU1WithTemp(CONVI64_TO_F)
     I64Shift(I64SHL)
     I64Shift(I64SHR)
     I64Shift(I64ASR)
@@ -2362,6 +2363,9 @@ namespace gbe
         sel.MOV(dst, unpacked);
       } else if ((dstType == ir::TYPE_S32 || dstType == ir::TYPE_U32) && srcFamily == FAMILY_QWORD) {
         sel.CONVI64_TO_I(dst, src);
+      } else if (dstType == ir::TYPE_FLOAT && srcFamily == FAMILY_QWORD) {
+        GenRegister tmp = sel.selReg(sel.reg(FAMILY_QWORD));
+        sel.CONVI64_TO_F(dst, src, tmp);
       } else if (dst.isdf()) {
         ir::Register r = sel.reg(ir::RegisterFamily::FAMILY_QWORD);
         sel.MOV_DF(dst, src, sel.selReg(r));
diff --git a/backend/src/backend/gen_insn_selection.hxx b/backend/src/backend/gen_insn_selection.hxx
index 5c857dd..a5e26fd 100644
--- a/backend/src/backend/gen_insn_selection.hxx
+++ b/backend/src/backend/gen_insn_selection.hxx
@@ -66,3 +66,4 @@ DECL_SELECTION_IR(UPSAMPLE_INT, BinaryInstruction)
 DECL_SELECTION_IR(UPSAMPLE_LONG, BinaryInstruction)
 DECL_SELECTION_IR(CONVI_TO_I64, UnaryWithTempInstruction)
 DECL_SELECTION_IR(CONVI64_TO_I, UnaryInstruction)
+DECL_SELECTION_IR(CONVI64_TO_F, UnaryWithTempInstruction)
diff --git a/kernels/compiler_long_convert.cl b/kernels/compiler_long_convert.cl
index 03df147..e5f7939 100644
--- a/kernels/compiler_long_convert.cl
+++ b/kernels/compiler_long_convert.cl
@@ -12,3 +12,8 @@ kernel void compiler_long_convert_2(global char *dst1, global short *dst2, globa
   dst2[i] = src[i];
   dst3[i] = src[i];
 }
+
+kernel void compiler_long_convert_to_float(global float *dst, global long *src) {
+  int i = get_global_id(0);
+  dst[i] = src[i];
+}
diff --git a/utests/compiler_long_convert.cpp b/utests/compiler_long_convert.cpp
index fe976be..8513f6d 100644
--- a/utests/compiler_long_convert.cpp
+++ b/utests/compiler_long_convert.cpp
@@ -116,3 +116,44 @@ void compiler_long_convert_2(void)
 }
 
 MAKE_UTEST_FROM_FUNCTION(compiler_long_convert_2);
+
+// convert 64-bit integer to 32-bit float
+void compiler_long_convert_to_float(void)
+{
+  const size_t n = 16;
+  int64_t src[n];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_convert", "compiler_long_convert_to_float");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int64_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  // Run random tests
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    src[i] = (int64_t)0x20000000 * i;
+  }
+  OCL_MAP_BUFFER(1);
+  memcpy(buf_data[1], src, sizeof(src));
+  OCL_UNMAP_BUFFER(1);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  float *dst = ((float *)buf_data[0]);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    //printf("%f\n", dst[i]);
+    OCL_ASSERT(dst[i] == src[i]);
+  }
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_long_convert_to_float);
+
-- 
1.8.1.2