Made test_imm_paramters / test_2d_copy pass on IVB

diff --git a/kernels/test_imm_parameters_kernels.cl b/kernels/test_imm_parameters_kernels.cl
index 433e8cc..dabcb69 100755 (executable)
--- a/kernels/test_imm_parameters_kernels.cl
+++ b/kernels/test_imm_parameters_kernels.cl
@@ -8,11 +8,12 @@ __kernel void test_imm_parameters (__global int *dst,
                                    const    int a5)
 {
   int tid = get_global_id(0);
-  dst[6*tid+0] = (int) a0 + src[tid];
-  dst[6*tid+1] = (int) a1 + src[tid];
-  dst[6*tid+2] = (int) a2 + src[tid];
-  dst[6*tid+3] = (int) a3 + src[tid];
-  dst[6*tid+4] = (int) a4 + src[tid];
-  dst[6*tid+5] = (int) a5 + src[tid];
+  const int from = src[tid];
+  dst[6*tid+0] = (int) a0 + from;
+  dst[6*tid+1] = (int) a1 + from;
+  dst[6*tid+2] = (int) a2 + from;
+  dst[6*tid+3] = (int) a3 + from;
+  dst[6*tid+4] = (int) a4 + from;
+  dst[6*tid+5] = (int) a5 + from;
 }
 

diff --git a/src/cl_kernel.c b/src/cl_kernel.c
index aece675..1efb159 100644 (file)
--- a/src/cl_kernel.c
+++ b/src/cl_kernel.c
@@ -349,7 +349,7 @@ cl_kernel_setup_patch_list(cl_kernel k, const char *patch, size_t sz)
       ASSOC_ITEM (THREAD_PAYLOAD, thread_payload, thread_payload);
 
       case PATCH_TOKEN_DATA_PARAMETER_STREAM:
-        info->curbe.sz = *(uint32_t *) patch;
+        info->curbe.sz = *(uint32_t *) (patch + sizeof(cl_patch_item_header_t));
         info->curbe.offset = 0;
       break;
       case PATCH_TOKEN_CONSTANT_MEMORY_KERNEL_ARGUMENT: