diff options
| author | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-05-31 10:09:56 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-05-31 15:21:39 +0800 |
| commit | 9c2f56daae96f51e39d4f504d84e5ef59bea77b5 (patch) | |
| tree | b7df440daa6599ad9e0407a7955c7047e78a05e7 /kernels/compiler_vector_load_store.cl | |
| parent | 8866634b02d4230c723c3c9449341ee749a0e2a4 (diff) | |
GBE: Fixed a 3 elements vector load/store bug.
Per OpenCL spec, for 3-component vector data types,the
size of the data type is 4 * sizeof(component). And llvm
FE really cast a type3 data to type4 data for load/store
instruction, thus break our implementation. We need to
fixup it to the actual element size.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.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 | 10 |
1 files changed, 5 insertions, 5 deletions
diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl index b362412c..30f0e1e3 100644 --- a/kernels/compiler_vector_load_store.cl +++ b/kernels/compiler_vector_load_store.cl @@ -18,12 +18,12 @@ __kernel void test_##type ##n(__global type *pin, \ } #define TEST_ALL_TYPE(n) \ - TEST_TYPE(char,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(ushort,n)\ + TEST_TYPE(int,n) \ + TEST_TYPE(uint,n) \ TEST_TYPE(float,n) #if 0 @@ -34,7 +34,7 @@ __kernel void test_##type ##n(__global type *pin, \ #endif TEST_ALL_TYPE(2) -//TEST_ALL_TYPE(3) +TEST_ALL_TYPE(3) TEST_ALL_TYPE(4) TEST_ALL_TYPE(8) TEST_ALL_TYPE(16) |