LGTM, will push it latter. Thanks.
On Tue, Sep 24, 2013 at 11:11:21AM +0800, Homer Hsing wrote: > llvm phi node can have odd number of args. > this patch also contains a test case. > > Signed-off-by: Homer Hsing <[email protected]> > --- > backend/src/llvm/llvm_scalarize.cpp | 1 - > kernels/compiler_vector_inc.cl | 13 +++++++++++ > utests/CMakeLists.txt | 1 + > utests/compiler_vector_inc.cpp | 46 > +++++++++++++++++++++++++++++++++++++ > 4 files changed, 60 insertions(+), 1 deletion(-) > create mode 100644 kernels/compiler_vector_inc.cl > create mode 100644 utests/compiler_vector_inc.cpp > > diff --git a/backend/src/llvm/llvm_scalarize.cpp > b/backend/src/llvm/llvm_scalarize.cpp > index 41674b6..7a40616 100644 > --- a/backend/src/llvm/llvm_scalarize.cpp > +++ b/backend/src/llvm/llvm_scalarize.cpp > @@ -383,7 +383,6 @@ namespace gbe { > > if (PHINode* phi = dyn_cast<PHINode>(inst)) { > PHINode* res = PHINode::Create(GetBasicType(inst), > phi->getNumIncomingValues()); > - assert(args.size() % 2 == 0 && "Odd number of arguments for a PHI"); > > // Loop over pairs of operands: [Value*, BasicBlock*] > for (unsigned int i = 0; i < args.size(); i++) { > diff --git a/kernels/compiler_vector_inc.cl b/kernels/compiler_vector_inc.cl > new file mode 100644 > index 0000000..548dcb4 > --- /dev/null > +++ b/kernels/compiler_vector_inc.cl > @@ -0,0 +1,13 @@ > +kernel void compiler_vector_inc(global char *dst, global char *src) { > + size_t i = get_global_id(0); > + char2 dst2 = vload2(i, dst); > + if (src[i] == 0) > + dst2++; > + else if(src[i] == 1) > + ++dst2; > + else if(src[i] == 2) > + dst2--; > + else > + --dst2; > + vstore2(dst2, i, dst); > +} > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index 24322b6..7e3c3ee 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -102,6 +102,7 @@ set (utests_sources > compiler_get_image_info.cpp > compiler_vect_compare.cpp > compiler_vector_load_store.cpp > + compiler_vector_inc.cpp > compiler_cl_finish.cpp > get_cl_info.cpp > builtin_atan2.cpp > diff --git a/utests/compiler_vector_inc.cpp b/utests/compiler_vector_inc.cpp > new file mode 100644 > index 0000000..abc5408 > --- /dev/null > +++ b/utests/compiler_vector_inc.cpp > @@ -0,0 +1,46 @@ > +#include <cstdint> > +#include <cstring> > +#include <iostream> > +#include "utest_helper.hpp" > + > +void compiler_vector_inc(void) > +{ > + const int n = 64; > + char dst[n]; > + char src[n]; > + > + OCL_CREATE_KERNEL("compiler_vector_inc"); > + OCL_CREATE_BUFFER(buf[0], 0, n, NULL); > + OCL_CREATE_BUFFER(buf[1], 0, n, NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); > + globals[0] = n / 2; > + locals[0] = 16; > + > + for (int i = 0; i < n; ++i) { > + dst[i] = i; > + src[i] = (i / 2) % 4; > + } > + OCL_MAP_BUFFER(0); > + OCL_MAP_BUFFER(1); > + memcpy(buf_data[0], dst, n); > + memcpy(buf_data[1], src, n); > + OCL_UNMAP_BUFFER(0); > + OCL_UNMAP_BUFFER(1); > + > + OCL_NDRANGE(1); > + > + OCL_MAP_BUFFER(0); > + char *dest = ((char *)buf_data[0]); > + for (int i=0; i<n; ++i) { > + char wish; > + if (src[i/2] < 2) > + wish = dst[i] + 1; > + else > + wish = dst[i] - 1; > + OCL_ASSERT(dest[i] == wish); > + } > + OCL_UNMAP_BUFFER(0); > +} > + > +MAKE_UTEST_FROM_FUNCTION(compiler_vector_inc); > -- > 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
