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

[Zhigang Gong]

Please ignore the change in kernels/test_copy_image.cl. That is for testing
purpose and was added into this patch by accident.

-----Original Message-----
From: beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
[mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org]
On Behalf Of Zhigang Gong
Sent: Wednesday, November 13, 2013 8:35 AM
To: beignet at lists.freedesktop.org
Cc: Zhigang Gong
Subject: [Beignet] [PATCH 1/2] GBE: remove all vstore macros for constant
memory space.

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

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet