Hi, If I change the line: src[i] = (int64_t)0x20000000 * i; to: src[i] = (int64_t)(-i); in the test case, it will fail.
I think the algorithm you change the long to float don't cover all case. For example change singed long -1L to float. -----Original Message----- From: [email protected] [mailto:[email protected]] On Behalf Of Homer Hsing Sent: Monday, September 02, 2013 10:12 AM To: [email protected] Subject: [Beignet] [PATCH] support converting 64-bit integer to 32-bit float Signed-off-by: Homer Hsing <[email protected]> --- 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 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
