[Beignet] [PATCH 2/2] utests: add more constant test cases for composite type.

Ruiling Song ruiling.song at intel.com
Tue Sep 17 19:18:43 PDT 2013


Signed-off-by: Ruiling Song <ruiling.song at intel.com>
---
 kernels/compiler_global_constant.cl |   59 +++++++++++++++++++++++++++-
 utests/compiler_global_constant.cpp |   74 +++++++++++++++++++++++++++++++++++
 2 files changed, 131 insertions(+), 2 deletions(-)

diff --git a/kernels/compiler_global_constant.cl b/kernels/compiler_global_constant.cl
index 5db58d6..71fe86c 100644
--- a/kernels/compiler_global_constant.cl
+++ b/kernels/compiler_global_constant.cl
@@ -1,10 +1,65 @@
 constant int m[3] = {71,72,73};
 constant int n = 1;
-constant int o[3] = {1, 1, 1};
+constant int o[3] = {3, 2, 1};
+
+constant int4 a= {1, 2, 3, 4};
+constant int4 b = {0, -1, -2, -3};
+
+struct Person {
+  char name[7];
+  int3 idNumber;
+};
+
+struct Test1 {
+  int a0;
+  char a1;
+};
+
+struct Test2 {
+  char a0;
+  int a1;
+};
+
+constant struct Person james= {{"james"}, (int3)(1, 2, 3)};
+
+constant struct Test1 t0 = {1, 2};
+constant struct Test2 t1 = {1, 2};
+
+constant int3 c[3] = {(int3)(0, 1, 2), (int3)(3, 4, 5), (int3)(6,7,8) };
+constant char4 d[3] = {(char4)(0, 1, 2, 3), (char4)(4, 5, 6, 7), (char4)(8, 9, 10, 11)};
+
+constant struct Person members[3] = {{{"abc"}, (int3)(1, 2, 3)}, { {"defg"}, (int3)(4,5,6)}, { {"hijk"}, (int3)(7,8,9)} };
 
 __kernel void
 compiler_global_constant(__global int *dst, int e, int r)
 {
   int id = (int)get_global_id(0);
-  dst[id] = m[id%3] * n * o[2] + e + r;
+
+  int4 x = a + b;
+  dst[id] = m[id%3] * n * o[2] + e + r *x.y * a.x;
+}
+// array of vectors
+__kernel void
+compiler_global_constant1(__global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = c[id%3].y + d[id%3].w;
+}
+
+// structure
+__kernel void
+compiler_global_constant2(__global int *dst)
+{
+  int id = (int)get_global_id(0);
+
+  dst[id] = james.idNumber.y + t0.a1 + t1.a1;
+}
+
+//array of structure
+__kernel void
+compiler_global_constant3(__global int *dst)
+{
+  int id = (int)get_global_id(0);
+
+  dst[id] = members[id%3].idNumber.z + members[id%3].name[2];
 }
diff --git a/utests/compiler_global_constant.cpp b/utests/compiler_global_constant.cpp
index c0bee4d..a2d0172 100644
--- a/utests/compiler_global_constant.cpp
+++ b/utests/compiler_global_constant.cpp
@@ -27,4 +27,78 @@ void compiler_global_constant(void)
   OCL_UNMAP_BUFFER(0);
 }
 
+void compiler_global_constant1(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_global_constant", "compiler_global_constant1");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  uint32_t data1[] = {1, 4, 7};
+  uint32_t data2[]= {3, 7, 11};
+
+  // Check results
+  OCL_MAP_BUFFER(0);
+  for (uint32_t i = 0; i < n; ++i)
+//    printf("%d result %d reference %d\n", i, ((uint32_t *)buf_data[0])[i], data1[i%3] + data2[i%3]);
+    OCL_ASSERT(((uint32_t *)buf_data[0])[i] == data1[i%3] + data2[i%3]);
+  OCL_UNMAP_BUFFER(0);
+}
+
+void compiler_global_constant2(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_global_constant", "compiler_global_constant2");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check results
+  OCL_MAP_BUFFER(0);
+  for (uint32_t i = 0; i < n; ++i)
+//    printf("%d result %d reference %d\n", i, ((uint32_t *)buf_data[0])[i], 6);
+    OCL_ASSERT(((uint32_t *)buf_data[0])[i] == 6);
+  OCL_UNMAP_BUFFER(0);
+}
+
+void compiler_global_constant3(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_global_constant", "compiler_global_constant3");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  uint32_t data1[] = {3, 6, 9};
+  char data2[]= {'c', 'f', 'j'};
+  // Check results
+  OCL_MAP_BUFFER(0);
+  for (uint32_t i = 0; i < n; ++i)
+//    printf("%d result %d reference %d\n", i, ((uint32_t *)buf_data[0])[i], data1[i%3] + (int)data2[i%3]);
+    OCL_ASSERT(((uint32_t *)buf_data[0])[i] == data1[i%3] + (uint32_t)data2[i%3]);
+  OCL_UNMAP_BUFFER(0);
+}
+
 MAKE_UTEST_FROM_FUNCTION(compiler_global_constant);
+MAKE_UTEST_FROM_FUNCTION(compiler_global_constant1);
+MAKE_UTEST_FROM_FUNCTION(compiler_global_constant2);
+MAKE_UTEST_FROM_FUNCTION(compiler_global_constant3);
-- 
1.7.9.5



More information about the Beignet mailing list