tensor test: Fix OpenCL kernel

src/datawizard/interfaces/tensor_interface.c

@@ -682,7 +682,6 @@ static int copy_opencl_common(void *src_interface, unsigned src_node, void *dst_
 	if (src_block->ldy == dst_block->ldy && src_block->ldz == dst_block->ldz && src_block->ldt == dst_block->ldt
 		&& nx*ny*nz == src_block->ldt)
 	{
-		/* FIXME: tests/datawizard/interfaces/tensor/tensor_interface reports this as bogus !? */
 		ret = starpu_opencl_copy_async_sync(src_block->dev_handle, src_block->offset, src_node,
 						    dst_block->dev_handle, dst_block->offset, dst_node,
 						    src_block->nx*src_block->ny*src_block->nz*src_block->nt*src_block->elemsize,

tests/datawizard/interfaces/tensor/tensor_interface.c

@@ -151,8 +151,7 @@ main(int argc, char **argv)
 	struct starpu_conf conf;
 	starpu_conf_init(&conf);
 	conf.ncuda = 2;
-	// FIXME: bogus for now !?
-	conf.nopencl = 0;
+	conf.nopencl = 1;
 	conf.nmic = -1;
 
 	if (starpu_initialize(&conf, &argc, &argv) == -ENODEV || starpu_cpu_worker_get_count() == 0)

tests/datawizard/interfaces/tensor/tensor_opencl_kernel.cl

@@ -33,14 +33,14 @@ __kernel void tensor_opencl(__global int *tensor,
 		{
 			for (i = 0; i < nx; i++)
 			{
-                                if (tensor[(k*ldz)+(j*ldy)+i] != factor * val)
+                                if (tensor[(l*ldt)+(k*ldz)+(j*ldy)+i] != factor * val)
 				{
 					*err = 1;
 					return;
 				}
 				else
 				{
-					tensor[(k*ldz)+(j*ldy)+i] *= -1;
+					tensor[(l*ldt)+(k*ldz)+(j*ldy)+i] *= -1;
 					val++;
 				}
 			}