diff options
author | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-05-24 17:42:18 +0800 |
---|---|---|
committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-05-29 13:58:27 +0800 |
commit | f289075a6928e8b2f5243dbca2c0df2e5628f582 (patch) | |
tree | bda1105bf0f742ea514c1b36e2cdcb5bc1e95f1b /kernels/compiler_vector_load_store.cl | |
parent | f017fe759939edd75ba09288327b81120d6805a3 (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.cl | 52 |
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) |