The patchset is pushed, with mirror changes. GBE: a potential bug in instruction scheduling. + || node->insn.opcode == SEL_OP_ELSE
utests: refine image 1d buffer test case. - int offset = x; > -----Original Message----- > From: Beignet [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of > Zhigang Gong > Sent: Thursday, August 13, 2015 10:24 > To: beignet@lists.freedesktop.org > Cc: Gong, Zhigang > Subject: [Beignet] [PATCH 2/4] utests: refine image 1d buffer test case. > > We need to test large image 1d buffer read and write testing. > > Signed-off-by: Zhigang Gong <zhigang.g...@intel.com> > --- > kernels/image_1D_buffer.cl | 11 ++----- utests/image_1D_buffer.cpp | 73 > ++++++++++++++++++---------------------------- > 2 files changed, 32 insertions(+), 52 deletions(-) > > diff --git a/kernels/image_1D_buffer.cl b/kernels/image_1D_buffer.cl index > e8e0a86..4f982e6 100644 > --- a/kernels/image_1D_buffer.cl > +++ b/kernels/image_1D_buffer.cl > @@ -1,13 +1,8 @@ > -__kernel void image_1D_buffer(image1d_buffer_t image1, image1d_t > image2, sampler_t sampler, __global int *results) > +__kernel void image_1D_buffer(image1d_buffer_t image1, > image1d_buffer_t > +image2) > { > int x = get_global_id(0); > int offset = x; > > - int4 col = read_imagei(image1, x); > - int4 test = (col != read_imagei(image2, sampler, x)); > - > - if (test.x || test.y || test.z || test.w) > - results[offset] = 0; > - else > - results[offset] = 1; > + uint4 color = read_imageui(image1, x); > + write_imageui(image2, x, color); > } > diff --git a/utests/image_1D_buffer.cpp b/utests/image_1D_buffer.cpp > index d8d761f..e2cfcde 100644 > --- a/utests/image_1D_buffer.cpp > +++ b/utests/image_1D_buffer.cpp > @@ -3,78 +3,63 @@ > > void image_1D_buffer(void) > { > - size_t buffer_sz = 1024; > - char *buf_content = (char *)malloc(buffer_sz * sizeof(char)); > + size_t buffer_sz = 65536; > + char *buf_content = (char *)malloc(buffer_sz * sizeof(int)); > int error; > cl_image_desc image_desc; > cl_image_format image_format; > - cl_sampler sampler; > - cl_mem image1, image2; > cl_mem ret_mem = NULL; > > OCL_CREATE_KERNEL("image_1D_buffer"); > > for (int32_t i = 0; i < (int32_t)buffer_sz; ++i) > - buf_content[i] = (rand() & 127); > + buf_content[i] = (rand() & 0xFFFFFFFF); > > - cl_mem buff = clCreateBuffer(ctx, CL_MEM_READ_ONLY | > CL_MEM_COPY_HOST_PTR, > - buffer_sz, buf_content, &error); > - OCL_ASSERT(error == CL_SUCCESS); > + OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_ONLY | > CL_MEM_COPY_HOST_PTR, > + buffer_sz * sizeof(int), buf_content); OCL_CREATE_BUFFER(buf[1], > + CL_MEM_READ_WRITE, buffer_sz * sizeof(int), NULL); > > memset(&image_desc, 0x0, sizeof(cl_image_desc)); > memset(&image_format, 0x0, sizeof(cl_image_format)); > > image_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; > - image_desc.image_row_pitch = buffer_sz; > - image_desc.image_width = buffer_sz / sizeof(uint32_t); //assume rgba32 > - image_desc.buffer = buff; > - > - image_format.image_channel_order = CL_RGBA; > - image_format.image_channel_data_type = CL_UNSIGNED_INT8; > - > - image1 = clCreateImage(ctx, CL_MEM_READ_ONLY, &image_format, > - &image_desc, NULL, &error ); > - OCL_ASSERT(error == CL_SUCCESS); > - > - error = clGetImageInfo(image1, CL_IMAGE_BUFFER, sizeof(ret_mem), > &ret_mem, NULL); > - OCL_ASSERT(error == CL_SUCCESS); > - OCL_ASSERT(ret_mem == buff); > + image_desc.image_row_pitch = buffer_sz * sizeof(int); > + image_desc.image_width = buffer_sz; //assume r32 image_desc.buffer = > + buf[0]; > > + image_format.image_channel_order = CL_R; > + image_format.image_channel_data_type = CL_UNSIGNED_INT32; > > - memset(&image_desc, 0x0, sizeof(cl_image_desc)); > - image_desc.image_type = CL_MEM_OBJECT_IMAGE1D; > - image_desc.image_width = buffer_sz / sizeof(uint32_t); > - image2 = clCreateImage(ctx, > CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, > - &image_format, &image_desc, buf_content, &error); > + // Create the source image1d_buffer. > + OCL_CREATE_IMAGE(buf[2], CL_MEM_READ_ONLY, &image_format, > + &image_desc, NULL); error = clGetImageInfo(buf[2], CL_IMAGE_BUFFER, > + sizeof(ret_mem), &ret_mem, NULL); > OCL_ASSERT(error == CL_SUCCESS); > + OCL_ASSERT(ret_mem == buf[0]); > > - // Create sampler to use > - sampler = clCreateSampler(ctx, false, CL_ADDRESS_NONE, > CL_FILTER_NEAREST, &error ); > + // Create the destination image1d_buffer. > + image_desc.buffer = buf[1]; > + OCL_CREATE_IMAGE(buf[3], CL_MEM_READ_ONLY, &image_format, > + &image_desc, NULL); error = clGetImageInfo(buf[3], CL_IMAGE_BUFFER, > + sizeof(ret_mem), &ret_mem, NULL); > OCL_ASSERT(error == CL_SUCCESS); > + OCL_ASSERT(ret_mem == buf[1]); > > - cl_mem result_buf = buf[0] = clCreateBuffer(ctx, 0, buffer_sz, NULL, > &error); > - OCL_ASSERT(error == CL_SUCCESS); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[2]); OCL_SET_ARG(1, > + sizeof(cl_mem), &buf[3]); > > - OCL_SET_ARG(0, sizeof(cl_mem), &image1); > - OCL_SET_ARG(1, sizeof(cl_mem), &image2); > - OCL_SET_ARG(2, sizeof(sampler), &sampler); > - OCL_SET_ARG(3, sizeof(cl_mem), &result_buf); > - > - globals[0] = buffer_sz/sizeof(int32_t); > + globals[0] = buffer_sz; > locals[0] = 16; > > OCL_NDRANGE(1); > > /* Now check the result. */ > OCL_MAP_BUFFER(0); > - for (uint32_t i = 0; i < buffer_sz/sizeof(int32_t); i++) > - OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 1); > + OCL_MAP_BUFFER(1); > + for (uint32_t i = 0; i < buffer_sz; i++) { > + if (((uint32_t*)buf_data[1])[i] != ((uint32_t*)buf_data[0])[i]) > + printf("i %d expected %x got %x \n", i, ((uint32_t*)buf_data[0])[i], > ((uint32_t*)buf_data[1])[i]); > + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == > + ((uint32_t*)buf_data[0])[i]); } > OCL_UNMAP_BUFFER(0); > - > - clReleaseSampler(sampler); > - clReleaseMemObject(image1); > - clReleaseMemObject(image2); > - clReleaseMemObject(buff); > + OCL_UNMAP_BUFFER(1); > } > > MAKE_UTEST_FROM_FUNCTION(image_1D_buffer); > -- > 1.9.1 > > _______________________________________________ > Beignet mailing list > Beignet@lists.freedesktop.org > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet
