tests/datawizard/scratch.c: Fix the OpenCL codelet.

tests/datawizard/scratch_opencl.c

@@ -29,6 +29,7 @@ void opencl_f(void *buffers[], void *args)
 	cl_event event;
 
 	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	unsigned elemsize = STARPU_VECTOR_GET_ELEMSIZE(buffers[0]);
 	cl_mem val = (cl_mem) STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
 	cl_mem tmp = (cl_mem) STARPU_VECTOR_GET_DEV_HANDLE(buffers[1]);
 
@@ -39,6 +40,18 @@ void opencl_f(void *buffers[], void *args)
 	if (err != CL_SUCCESS)
 		STARPU_OPENCL_REPORT_ERROR(err);
 
+	err = clEnqueueCopyBuffer(queue,
+		val,
+		tmp,
+		0,           /* offset in val */
+		0,           /* offset in tmp */
+		n * elemsize,
+		0,           /* num_events_in_wait_list */
+		NULL,        /* event_wait_list */
+		NULL);       /* event */
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
 	err = clSetKernelArg(kernel, 0, sizeof(val), &val);
 	err|= clSetKernelArg(kernel, 1, sizeof(tmp), &tmp);
 	err|= clSetKernelArg(kernel, 2, sizeof(n), &n);

tests/datawizard/scratch_opencl_kernel.cl

@@ -18,12 +18,8 @@ __kernel void increment_vector_opencl(__global unsigned *val,
 				      __global unsigned *tmp,
 				      unsigned nx)
 {
-        const int tid = get_global_id(0);
-	const uint nthreads = get_local_size(0);
+        const int id = get_global_id(0);
 
-	int i;
-	for (i = tid; i < nx; i += nthreads)
-	{
-		val[i] = tmp[i] + 1;
-	}
+	if (id < nx)
+		val[id] = tmp[id] + 1;
 }