GBE: Refine bti usage in backend & runtime.

Previously, we simply map 2G surface for memory access,
which has obvious security issue, user can easily read/write graphics
memory that does not belong to him. To prevent such kind of behaviour,
We bind each surface to a dedicated bti. HW provides automatic
bounds check. For out-of-bound write, it will be ignored. And for read
out-of-bound, hardware will simply return zero value.

The idea behind the patch is for a load/store instruction, it will search
through the LLVM use-def chain until finding out where the address
comes from. Then the bti is saved in ir::Instruction and used for
the later code generation. And for mixed pointer case, a load/store
will access more than one bti.

To simplify some code, '0' is reserved for constant address space,
'1' is reserved for private address space. Other btis are assigned
automatically by backend.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

1 files changed, 23 insertions, 0 deletions

diff --git a/kernels/compiler_mixed_pointer.cl b/kernels/compiler_mixed_pointer.cl
new file mode 100644
index 00000000..78c57835
--- /dev/null
+++ b/kernels/compiler_mixed_pointer.cl
@@ -0,0 +1,23 @@
+
+kernel void compiler_mixed_pointer(__global uint* src1, __global uint *src2, __global uint *dst) {
+  int x = get_global_id(0);
+  global uint * tmp = NULL;
+
+  switch(x) {
+    case 0:
+    case 1:
+    case 4:
+      tmp = src1;
+      break;
+    default:
+      tmp = src2;
+      break;
+  }
+  dst[x] = tmp[x];
+}
+
+kernel void compiler_mixed_pointer1(__global uint* src, __global uint *dst1, __global uint *dst2) {
+  int x = get_global_id(0);
+  global uint * tmp = x < 5 ? dst1 : dst2;
+  tmp[x] = src[x];
+}