summaryrefslogtreecommitdiff
path: root/kernels/compiler_vector_load_store.cl
diff options
context:
space:
mode:
authorZhigang Gong <zhigang.gong@linux.intel.com>2013-05-24 17:42:18 +0800
committerZhigang Gong <zhigang.gong@linux.intel.com>2013-05-29 13:58:27 +0800
commitf289075a6928e8b2f5243dbca2c0df2e5628f582 (patch)
treebda1105bf0f742ea514c1b36e2cdcb5bc1e95f1b /kernels/compiler_vector_load_store.cl
parentf017fe759939edd75ba09288327b81120d6805a3 (diff)
utests: test vector load and store.
Add float4/short4/char4 test case. Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com> Tested-by: Lv, Meng <meng.lv@intel.com> Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Diffstat (limited to 'kernels/compiler_vector_load_store.cl')
-rw-r--r--kernels/compiler_vector_load_store.cl52
1 files changed, 37 insertions, 15 deletions
diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl
index 28fd93a8..b362412c 100644
--- a/kernels/compiler_vector_load_store.cl
+++ b/kernels/compiler_vector_load_store.cl
@@ -1,18 +1,40 @@
/* test OpenCL 1.1 Vector Data Load/Store Functions (section 6.11.7) */
-kernel void compiler_vector_load_store() {
- float p[16], f;
- float4 f4;
- f4 = vload4(0, p);
- vstore4(f4, 0, p);
-
- long x[16], l;
- long16 l16;
- l = vload16(0, x);
- vstore16(l16, 0, x);
- half h[16];
- half4 h4;
- f = vload_half(0, h);
- f4 = vload_half4(0, h);
- vstore_half(f, 0, h);
+#define OFFSET2(type) (type ##2) {(type)1, (type)2}
+#define OFFSET3(type) (type ##3) {(type)1, (type)2, (type)3}
+#define OFFSET4(type) (type ##4) {(type)1, (type)2, (type)3, (type)4}
+#define OFFSET8(type) (type ##8) {(type)1, (type)2, (type)3, (type)4, (type)5, (type)6, (type)7, (type)8}
+#define OFFSET16(type) (type ##16) {(type)1, (type)2, (type)3, (type)4, (type)5, (type)6, (type)7, (type)8, (type)9, (type)10, (type)11, (type)12, (type)13, (type)14, (type)15, (type)16}
+
+#define TEST_TYPE(type, n) \
+__kernel void test_##type ##n(__global type *pin, \
+ __global type *pout) \
+{\
+ int x = get_global_id(0); \
+ type ##n value; \
+ value = vload ##n(x, pin); \
+ value += OFFSET ##n(type); \
+ vstore ##n(value, x, pout); \
}
+
+#define TEST_ALL_TYPE(n) \
+ TEST_TYPE(char,n) \
+ TEST_TYPE(uchar,n) \
+ TEST_TYPE(short,n) \
+ TEST_TYPE(ushort,n) \
+ TEST_TYPE(int,n) \
+ TEST_TYPE(uint,n) \
+ TEST_TYPE(float,n)
+
+#if 0
+ TEST_TYPE(double,n)
+ TEST_TYPE(long,n)
+ TEST_TYPE(ulong,n)
+ TEST_TYPE(half,n)
+#endif
+
+TEST_ALL_TYPE(2)
+//TEST_ALL_TYPE(3)
+TEST_ALL_TYPE(4)
+TEST_ALL_TYPE(8)
+TEST_ALL_TYPE(16)