[Beignet] [PATCH 1/2] GBE: remove all vstore macros for constant memory space.

[Zhigang Gong]

Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
---
 backend/src/ocl_stdlib.tmpl.h |   20 +++++++++++++++++++-
 kernels/test_copy_image.cl    |    8 +++++---
 2 files changed, 24 insertions(+), 4 deletions(-)

diff --git a/backend/src/ocl_stdlib.tmpl.h b/backend/src/ocl_stdlib.tmpl.h
index 7af39ce..23c8f6a 100644
--- a/backend/src/ocl_stdlib.tmpl.h
+++ b/backend/src/ocl_stdlib.tmpl.h
@@ -1860,6 +1860,11 @@ INLINE_OVERLOADABLE void vstore##DIM(TYPE##DIM v, size_t offset, SPACE TYPE *p)
   *(SPACE TYPE##DIM *) (p + DIM * offset) = v; \
 }
 
+#define DECL_UNTYPED_RD_SPACE_N(TYPE, DIM, SPACE) \
+INLINE_OVERLOADABLE TYPE##DIM vload##DIM(size_t offset, const SPACE TYPE *p) { \
+  return *(SPACE TYPE##DIM *) (p + DIM * offset); \
+}
+
 #define DECL_UNTYPED_V3_SPACE(TYPE, SPACE) \
 INLINE_OVERLOADABLE void vstore3(TYPE##3 v, size_t offset, SPACE TYPE *p) {\
   *(p + 3 * offset) = v.s0; \
@@ -1870,6 +1875,12 @@ INLINE_OVERLOADABLE TYPE##3 vload3(size_t offset, const SPACE TYPE *p) { \
   return *(SPACE TYPE##3 *) (p + 3 * offset); \
 }
 
+#define DECL_UNTYPED_RDV3_SPACE(TYPE, SPACE) \
+INLINE_OVERLOADABLE TYPE##3 vload3(size_t offset, const SPACE TYPE *p) { \
+  return *(SPACE TYPE##3 *) (p + 3 * offset); \
+}
+
+
 #define DECL_UNTYPED_RW_ALL_SPACE(TYPE, SPACE) \
   DECL_UNTYPED_RW_SPACE_N(TYPE, 2, SPACE) \
   DECL_UNTYPED_V3_SPACE(TYPE, SPACE) \
@@ -1877,10 +1888,17 @@ INLINE_OVERLOADABLE TYPE##3 vload3(size_t offset, const SPACE TYPE *p) { \
   DECL_UNTYPED_RW_SPACE_N(TYPE, 8, SPACE) \
   DECL_UNTYPED_RW_SPACE_N(TYPE, 16, SPACE)
 
+#define DECL_UNTYPED_RD_ALL_SPACE(TYPE, SPACE) \
+  DECL_UNTYPED_RD_SPACE_N(TYPE, 2, SPACE) \
+  DECL_UNTYPED_RDV3_SPACE(TYPE, SPACE) \
+  DECL_UNTYPED_RD_SPACE_N(TYPE, 4, SPACE) \
+  DECL_UNTYPED_RD_SPACE_N(TYPE, 8, SPACE) \
+  DECL_UNTYPED_RD_SPACE_N(TYPE, 16, SPACE)
+
 #define DECL_UNTYPED_RW_ALL(TYPE) \
   DECL_UNTYPED_RW_ALL_SPACE(TYPE, __global) \
   DECL_UNTYPED_RW_ALL_SPACE(TYPE, __local) \
-  DECL_UNTYPED_RW_ALL_SPACE(TYPE, __constant) \
+  DECL_UNTYPED_RD_ALL_SPACE(TYPE, __constant) \
   DECL_UNTYPED_RW_ALL_SPACE(TYPE, __private)
 
 DECL_UNTYPED_RW_ALL(char)
diff --git a/kernels/test_copy_image.cl b/kernels/test_copy_image.cl
index a5ee5e8..e0b586e 100644
--- a/kernels/test_copy_image.cl
+++ b/kernels/test_copy_image.cl
@@ -1,10 +1,12 @@
 __kernel void
-test_copy_image(__read_only image2d_t src, __write_only image2d_t dst, sampler_t sampler)
+test_copy_image(__read_only image2d_t src, __write_only image2d_t dst, sampler_t sampler, __global char*buff)
 {
   int2 coord;
   int4 color;
+  constant char* foo = "" __FILE__;
   coord.x = (int)get_global_id(0);
   coord.y = (int)get_global_id(1);
-  color = read_imagei(src, sampler, coord);
-  write_imagei(dst, coord, color);
+  //color = read_imagei(src, sampler, coord);
+ // write_imagei(dst, coord, color);
+  buff[get_global_id(0)] = foo[get_global_id(0)];
 }
-- 
1.7.9.5