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 98db273..b0ac53a 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -136,6 +136,7 @@ set (utests_sources
compiler_long_mult.cpp
compiler_long_cmp.cpp
compiler_bool_cross_basic_block.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
