[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