Hi chuanbo,
I just tried this case on my machine. It just passed as below:
compiler_vstore3_no_workgroup:
compiler_vstore3_no_workgroup() [SUCCESS]
Did you use the latest master branch to do the test?
On Tue, Aug 27, 2013 at 10:47:37PM +0800, Chuanbo Weng wrote:
> When there is no workgroup, the vstore3 will store incorrect
> value to buffer.
>
> Signed-off-by: Chuanbo Weng <chuanbo.w...@intel.com>
> ---
> kernels/compiler_vstore3_no_workgroup.cl | 11 +++++++++
> utests/CMakeLists.txt | 1 +
> utests/compiler_vstore3_no_workgroup.cpp | 39
> ++++++++++++++++++++++++++++++
> 3 files changed, 51 insertions(+)
> create mode 100644 kernels/compiler_vstore3_no_workgroup.cl
> create mode 100644 utests/compiler_vstore3_no_workgroup.cpp
>
> diff --git a/kernels/compiler_vstore3_no_workgroup.cl
> b/kernels/compiler_vstore3_no_workgroup.cl
> new file mode 100644
> index 0000000..aa641ae
> --- /dev/null
> +++ b/kernels/compiler_vstore3_no_workgroup.cl
> @@ -0,0 +1,11 @@
> +__kernel
> +void compiler_vstore3_no_workgroup(__global uint *vectorArray){
> + int col = get_global_id(0);
> +
> + uint3 val3;
> + val3.x = col*3 + 0;
> + val3.y = col*3 + 1;
> + val3.z = col*3 + 2;
> + size_t thumb_offset = col;
> + vstore3(val3, thumb_offset, vectorArray);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 97b7519..d9cb024 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -135,6 +135,7 @@ set (utests_sources
> compiler_long_asr.cpp
> compiler_long_mult.cpp
> compiler_long_cmp.cpp
> + compiler_vstore3_no_workgroup.cpp
> utest_assert.cpp
> utest.cpp
> utest_file_map.cpp
> diff --git a/utests/compiler_vstore3_no_workgroup.cpp
> b/utests/compiler_vstore3_no_workgroup.cpp
> new file mode 100644
> index 0000000..96491cd
> --- /dev/null
> +++ b/utests/compiler_vstore3_no_workgroup.cpp
> @@ -0,0 +1,39 @@
> +#include "utest_helper.hpp"
> +
> +static void cpu(int global_id, unsigned int *dst) {
> + for(unsigned int j = 0; j < 3; j++){
> + *(dst + global_id*3 + j) = global_id*3 + j;
> + }
> +}
> +
> +void compiler_vstore3_no_workgroup(void){
> +
> + const size_t n = 16;
> + unsigned int cpu_dst[16*3];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("compiler_vstore3_no_workgroup");
> + OCL_CREATE_BUFFER(buf[0], 0, n*3*sizeof(unsigned int), NULL);
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> + globals[0] = n;
> +
> + // Run the kernel on GPU
> + OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, 1, NULL, globals, NULL,
> 0, NULL, NULL);
> +
> + // Run on CPU
> + for (int32_t i = 0; i < (int32_t) n; ++i)
> + cpu(i, cpu_dst);
> +
> + // Compare
> + OCL_MAP_BUFFER(0);
> + unsigned int *px = (unsigned int *)buf_data[0];
> + for (int32_t i = 0; i < (int32_t) n; ++i)
> + for (int32_t j = 0; j < 3; ++j){
> + OCL_ASSERT(*(px+i*3+j) == *(cpu_dst+i*3+j));
> + }
> + OCL_UNMAP_BUFFER(0);
> +
> +}
> +
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_vstore3_no_workgroup)
> --
> 1.7.9.5
>
> _______________________________________________
> 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
