merge

tests/datawizard/interfaces/multiformat/multiformat_conversion_codelets_cuda.cu

@@ -0,0 +1,51 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <starpu.h>
+#include <starpu_cuda.h>
+#include "multiformat_types.h"
+#include "../../../helper.h"
+
+static __global__ void cpu_to_cuda_cuda(struct point *src,
+	struct struct_of_arrays *dst, unsigned n)
+{
+        unsigned i =  blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (i < n)
+	{
+		dst->x[i] = src[i].x;
+		dst->y[i] = src[i].y;
+	}
+
+}
+
+extern "C" void cpu_to_cuda_cuda_func(void *buffers[], void *_args)
+{
+	FPRINTF(stderr, "Entering %s\n", __func__);
+	struct point *src;
+	struct struct_of_arrays *dst;
+
+	src = (struct point *) STARPU_MULTIFORMAT_GET_CPU_PTR(buffers[0]);
+	dst = (struct struct_of_arrays *) STARPU_MULTIFORMAT_GET_CUDA_PTR(buffers[0]);
+
+	int n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);
+
+	unsigned threads_per_block = 64;
+	unsigned nblocks = (n + threads_per_block-1) / threads_per_block;
+
+        cpu_to_cuda_cuda<<<nblocks,threads_per_block,2,starpu_cuda_get_local_stream()>>>(src, dst, n);
+
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/multiformat/multiformat_conversion_codelets_kernel.cl

@@ -0,0 +1,27 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include "multiformat_types.h"
+__kernel void cpu_to_opencl_opencl(__global struct point *src,
+				   __global struct struct_of_arrays *dst,
+				   unsigned int n)
+{
+	const unsigned int i = get_global_id(0);
+	if (i < n)
+	{
+		dst->x[i] = src[i].x;
+		dst->y[i] = src[i].y;
+	}
+}

tests/datawizard/interfaces/multiformat/multiformat_conversion_codelets_opencl.c

@@ -0,0 +1,111 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ * Copyright (C) 2012  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include <config.h>
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include "../../../helper.h"
+
+#define KERNEL_LOCATION "tests/datawizard/interfaces/multiformat/multiformat_conversion_codelets_kernel.cl"
+static struct starpu_opencl_program opencl_conversion_program;
+
+void cpu_to_opencl_opencl_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	(void) args;
+	int id, devid, ret;
+        cl_int err;
+	cl_kernel kernel;
+	cl_command_queue queue;
+	cl_event event;
+
+	unsigned n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);
+	cl_mem src = (cl_mem) STARPU_MULTIFORMAT_GET_CPU_PTR(buffers[0]);
+	cl_mem dst = (cl_mem) STARPU_MULTIFORMAT_GET_OPENCL_PTR(buffers[0]);
+
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+
+	ret = starpu_opencl_load_opencl_from_file(KERNEL_LOCATION,
+						  &opencl_conversion_program,
+						  NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_conversion_program,
+					"cpu_to_opencl_opencl",
+					devid);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 0, sizeof(src), &src);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 1, sizeof(dst), &dst);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 2, sizeof(n), &n);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	{
+		size_t global=n;
+		size_t local;
+                size_t s;
+                cl_device_id device;
+
+                starpu_opencl_get_device(devid, &device);
+
+                err = clGetKernelWorkGroupInfo (kernel,
+						device,
+						CL_KERNEL_WORK_GROUP_SIZE,
+						sizeof(local),
+						&local,
+						&s);
+                if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+
+                if (local > global)
+			local = global;
+
+		err = clEnqueueNDRangeKernel(queue,
+					kernel,
+					1,
+					NULL,
+					&global,
+					&local,
+					0,
+					NULL,
+					&event);
+
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+        starpu_opencl_unload_opencl(&opencl_conversion_program);
+}

tests/datawizard/interfaces/multiformat/multiformat_cuda.cu

@@ -0,0 +1,77 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <starpu.h>
+#include <starpu_cuda.h>
+#include "multiformat_types.h"
+#include "../test_interfaces.h"
+#include "../../../helper.h"
+
+extern struct test_config multiformat_config;
+
+static __global__ void multiformat_cuda(struct struct_of_arrays *soa, unsigned n,
+					int *err, int factor)
+{
+        unsigned i =  blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (i >= n)
+		return;
+
+	if (soa->x[i] != i * factor || soa->y[i] != i * factor)
+	{
+		*err = 1;
+	}
+	else
+	{
+		soa->x[i] = -soa->x[i];
+		soa->y[i] = -soa->y[i];
+	}
+}
+
+extern "C" void test_multiformat_cuda_func(void *buffers[], void *args)
+{
+	FPRINTF(stderr, "Entering %s\n", __func__);
+	int factor;
+	int *ret;
+	cudaError_t error;
+	unsigned int n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);
+	struct struct_of_arrays *soa;
+
+	soa = (struct struct_of_arrays *) STARPU_MULTIFORMAT_GET_CUDA_PTR(buffers[0]);
+	unsigned threads_per_block = 64;
+	unsigned nblocks = (n + threads_per_block-1) / threads_per_block;
+	factor = *(int *) args;
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	error = cudaMemcpy(ret,
+			   &multiformat_config.copy_failed,
+			   sizeof(int),
+			   cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+        multiformat_cuda<<<nblocks,threads_per_block,2,starpu_cuda_get_local_stream()>>>(soa, n, ret, factor);
+
+	error = cudaMemcpy(&multiformat_config.copy_failed,
+			   ret,
+			   sizeof(int),
+			   cudaMemcpyDeviceToHost);
+	
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/multiformat/multiformat_opencl.c

@@ -0,0 +1,134 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <config.h>
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include "../test_interfaces.h"
+
+#define KERNEL_LOCATION "tests/datawizard/interfaces/multiformat/multiformat_opencl_kernel.cl"
+
+extern struct test_config multiformat_config;
+static struct starpu_opencl_program multiformat_program;
+
+void test_multiformat_opencl_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	int id, devid, factor, ret;
+	unsigned int n;
+
+        cl_int             err;
+	cl_kernel          kernel;
+	cl_command_queue   queue;
+	cl_event           event;
+	cl_context         context;
+	cl_mem             val, fail;
+
+	ret = starpu_opencl_load_opencl_from_file(KERNEL_LOCATION,
+						  &multiformat_program,
+						  NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+
+	factor = *(int *)args;
+	n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);
+	val = (cl_mem)STARPU_MULTIFORMAT_GET_OPENCL_PTR(buffers[0]);
+
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+	starpu_opencl_get_context(devid, &context);
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&multiformat_program,
+					"multiformat_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	fail = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+		sizeof(int), &multiformat_config.copy_failed, &err);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	/* Setting args */
+	err  = clSetKernelArg(kernel, 0, sizeof(val), &val);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 1, sizeof(n), &n);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 2, sizeof(fail), &fail);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 3, sizeof(factor), &factor);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		size_t global=n;
+		size_t local;
+                size_t s;
+                cl_device_id device;
+
+                starpu_opencl_get_device(devid, &device);
+
+                err = clGetKernelWorkGroupInfo (kernel,
+						device,
+						CL_KERNEL_WORK_GROUP_SIZE,
+						sizeof(local),
+						&local,
+						&s);
+                if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+
+                if (local > global)
+			local = global;
+
+		err = clEnqueueNDRangeKernel(queue,
+					kernel,
+					1,
+					NULL,
+					&global,
+					&local,
+					0,
+					NULL,
+					&event);
+
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	err = clEnqueueReadBuffer(queue,
+				  fail,
+				  CL_TRUE,
+				  0, 
+				  sizeof(int),
+				  &multiformat_config.copy_failed,
+				  0,
+				  NULL,
+				  NULL);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+        starpu_opencl_unload_opencl(&multiformat_program);
+}

tests/datawizard/interfaces/multiformat/multiformat_opencl_kernel.cl

@@ -0,0 +1,36 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "multiformat_types.h"
+__kernel void multiformat_opencl(__global struct struct_of_arrays *soa,
+				 unsigned int nx,
+				 __global int *err,
+				 int factor)
+{
+        const int i = get_global_id(0);
+	if (i >= nx)
+		return;
+
+	if (soa->x[i] != i * factor || soa->y[i] != i * factor)
+	{
+		*err = i;
+	}
+	else
+	{
+		soa->x[i] = -soa->x[i];
+		soa->y[i] = -soa->y[i];
+	}
+}

tests/datawizard/interfaces/multiformat/multiformat_types.h

@@ -0,0 +1,33 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#ifndef TEST_MULTIFORMAT_TYPES_H
+#define TEST_MULTIFORMAT_TYPES_H
+
+#define N_ELEMENTS 2
+
+struct struct_of_arrays
+{
+	int x[N_ELEMENTS];
+	int y[N_ELEMENTS];
+};
+
+struct point
+{
+	int x, y;
+};
+
+
+#endif

tests/datawizard/interfaces/variable/variable_cuda.cu

@@ -0,0 +1,67 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <starpu.h>
+#include <starpu_cuda.h>
+#include "../test_interfaces.h"
+
+extern struct test_config variable_config;
+
+static __global__ void variable_cuda(int *val, int *err, int factor)
+{
+        unsigned i =  blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (i > 0)
+		return;
+
+	if (*val != 42 * factor)
+		*err = 1;
+	else
+		*val *= -1;
+}
+
+extern "C" void test_variable_cuda_func(void *buffers[], void *args)
+{
+	cudaError_t error;
+	int *ret;
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	error = cudaMemcpy(ret,
+			   &variable_config.copy_failed,
+			   sizeof(int),
+			   cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+        int *val = (int *)STARPU_VARIABLE_GET_PTR(buffers[0]);
+	int factor = *(int*) args;
+
+	unsigned threads_per_block = 64;
+	unsigned nblocks = 1;
+
+        variable_cuda<<<nblocks,threads_per_block,0,starpu_cuda_get_local_stream()>>>(val, ret, factor);
+	error = cudaMemcpy(&variable_config.copy_failed,
+			   ret,
+			   sizeof(int),
+			   cudaMemcpyDeviceToHost);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/variable/variable_interface.c

@@ -0,0 +1,119 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <config.h>
+#include <starpu.h>
+#include "../test_interfaces.h"
+
+static int variable;
+static int variable2;
+static starpu_data_handle_t variable_handle;
+static starpu_data_handle_t variable2_handle;
+
+/* Codelets */
+#ifdef STARPU_USE_CPU
+static void test_variable_cpu_func(void *buffers[], void *args);
+#endif
+#ifdef STARPU_USE_CUDA
+extern void test_variable_cuda_func(void *buffers[], void *args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_variable_opencl_func(void *buffers[], void *args);
+#endif
+
+struct test_config variable_config =
+{
+#ifdef STARPU_USE_CPU
+	.cpu_func     = test_variable_cpu_func,
+#endif
+#ifdef STARPU_USE_CUDA
+	.cuda_func    = test_variable_cuda_func,
+#endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_func  = test_variable_opencl_func,
+#endif
+	.handle       = &variable_handle,
+	.dummy_handle = &variable2_handle,
+	.copy_failed  = 0,
+	.name         = "variable_interface"
+};
+
+static void
+test_variable_cpu_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	int *val;
+	int factor;
+
+	val = (int *) STARPU_VARIABLE_GET_PTR(buffers[0]);
+	factor = *(int *)args;
+
+	if (*val != 42 * factor)
+		variable_config.copy_failed = 1;
+	else
+		*val *= -1;
+}
+
+
+static
+void register_data(void)
+{
+	variable = 42;
+	variable2 = 12;
+
+	starpu_variable_data_register(&variable_handle, 0,
+				      (uintptr_t) &variable, sizeof(variable));
+	starpu_variable_data_register(&variable2_handle, 0,
+				      (uintptr_t) &variable2, sizeof(variable2));
+}
+
+static
+void unregister_data(void)
+{
+	starpu_data_unregister(variable_handle);
+	starpu_data_unregister(variable2_handle);
+}
+
+int
+main(void)
+{
+	int ret;
+	data_interface_test_summary *summary;
+
+	struct starpu_conf conf =
+	{
+		.ncpus = -1,
+		.ncuda = 2,
+		.nopencl = 1
+	};
+
+	ret = starpu_init(&conf);
+	if (ret == -ENODEV) return STARPU_TEST_SKIPPED;
+
+	register_data();
+
+	summary = run_tests(&variable_config);
+	if (!summary)
+		exit(EXIT_FAILURE);
+
+	unregister_data();
+
+	starpu_shutdown();
+
+	data_interface_test_summary_print(stderr, summary);
+
+	return data_interface_test_summary_success(summary);
+}

tests/datawizard/interfaces/variable/variable_opencl.c

@@ -0,0 +1,128 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <config.h>
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include "../test_interfaces.h"
+
+#define KERNEL_LOCATION "tests/datawizard/interfaces/variable/variable_opencl_kernel.cl"
+
+extern struct test_config variable_config;
+static struct starpu_opencl_program opencl_program;
+
+void test_variable_opencl_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	int id, devid, ret;
+	int factor = *(int *) args;
+
+        cl_int             err;
+	cl_kernel          kernel;
+	cl_command_queue   queue;
+	cl_event           event;
+
+	ret = starpu_opencl_load_opencl_from_file(KERNEL_LOCATION, &opencl_program, NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+
+	cl_mem val = (cl_mem)STARPU_VARIABLE_GET_PTR(buffers[0]);
+
+	cl_context context;
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+	starpu_opencl_get_context(devid, &context);
+
+	cl_mem fail = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+		sizeof(int), &variable_config.copy_failed, &err);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_program,
+					"variable_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err  = clSetKernelArg(kernel, 0, sizeof(val), &val);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 1, sizeof(fail), &fail);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 2, sizeof(factor), &factor);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		size_t global = 1;
+		size_t local;
+                size_t s;
+                cl_device_id device;
+
+                starpu_opencl_get_device(devid, &device);
+
+                err = clGetKernelWorkGroupInfo (kernel,
+						device,
+						CL_KERNEL_WORK_GROUP_SIZE,
+						sizeof(local),
+						&local,
+						&s);
+                if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+
+                if (local > global)
+			local = global;
+
+		err = clEnqueueNDRangeKernel(queue,
+					kernel,
+					1,
+					NULL,
+					&global,
+					&local,
+					0,
+					NULL,
+					&event);
+
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	err = clEnqueueReadBuffer(queue,
+				  fail,
+				  CL_TRUE,
+				  0, 
+				  sizeof(int),
+				  &variable_config.copy_failed,
+				  0,
+				  NULL,
+				  NULL);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+        starpu_opencl_unload_opencl(&opencl_program);
+	return;
+}

tests/datawizard/interfaces/variable/variable_opencl_kernel.cl

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+__kernel void variable_opencl(__global int *val,
+			      __global int *err,
+			      int factor)
+{
+        const int i = get_global_id(0);
+        if (i > 0)
+		return;
+
+	if (*val != 42 * factor)
+		*err = 1;
+	else
+		*val *= -1;
+}
+

tests/datawizard/interfaces/vector/test_vector_cuda.cu

@@ -0,0 +1,67 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <starpu.h>
+#include <starpu_cuda.h>
+#include "../test_interfaces.h"
+
+extern struct test_config vector_config;
+
+__global__ void framework_cuda(int *val, unsigned n, int *err, int factor)
+{
+        unsigned i =  blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (i >= n)
+		return;
+
+	if (val[i] != i*factor)
+		*err = 1;
+	else
+		val[i] = -val[i];
+}
+
+extern "C" void test_vector_cuda_func(void *buffers[], void *args)
+{
+	cudaError_t error;
+	int *ret;
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+	{
+		fprintf(stderr, "cudaMalloc failed...\n");
+		return;
+	}
+
+	error = cudaMemcpy(ret, &vector_config.copy_failed, sizeof(int), cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		return;
+
+        unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+        int *val = (int *)STARPU_VECTOR_GET_PTR(buffers[0]);
+	int factor = *(int*) args;
+
+	unsigned threads_per_block = 64;
+	unsigned nblocks = (n + threads_per_block-1) / threads_per_block;
+
+        framework_cuda<<<nblocks,threads_per_block,0,starpu_cuda_get_local_stream()>>>(val, n, ret, factor);
+	error = cudaMemcpy(&vector_config.copy_failed, ret, sizeof(int), cudaMemcpyDeviceToHost);
+	if (error != cudaSuccess)
+	{
+		return;
+	}
+
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/vector/test_vector_interface.c

@@ -0,0 +1,132 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <config.h>
+#include <starpu.h>
+#include "../test_interfaces.h"
+#include "../../../helper.h"
+
+/* Prototypes */
+static void register_data(void);
+static void unregister_data(void);
+static void test_vector_cpu_func(void *buffers[], void *args);
+#ifdef STARPU_USE_CUDA
+extern void test_vector_cuda_func(void *buffers[], void *_args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_vector_opencl_func(void *buffers[], void *args);
+#endif
+
+
+static starpu_data_handle_t vector_handle;
+static starpu_data_handle_t vector2_handle;
+
+struct test_config vector_config =
+{
+	.cpu_func      = test_vector_cpu_func,
+#ifdef STARPU_USE_CUDA
+	.cuda_func     = test_vector_cuda_func,
+#endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_func   = test_vector_opencl_func,
+#endif
+	.handle        = &vector_handle,
+	.dummy_handle  = &vector2_handle,
+	.copy_failed   = 0,
+	.name          = "vector_interface"
+};
+
+#define VECTOR_SIZE 123
+static int vector[VECTOR_SIZE];
+static int vector2[VECTOR_SIZE];
+
+static void
+register_data(void)
+{
+	/* Initializing data */
+	int i;
+	for (i = 0; i < VECTOR_SIZE; i++)
+		vector[i] = i;
+
+	/* Registering data */
+	starpu_vector_data_register(&vector_handle,
+                                    0,
+                                    (uintptr_t)vector,
+				    VECTOR_SIZE,
+				    sizeof(int));
+	starpu_vector_data_register(&vector2_handle,
+                                    0,
+                                    (uintptr_t)vector2,
+				    VECTOR_SIZE,
+				    sizeof(int));
+}
+
+static void
+unregister_data(void)
+{
+	starpu_data_unregister(vector_handle);
+	starpu_data_unregister(vector2_handle);
+}
+
+static void test_vector_cpu_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	int *val = (int *) STARPU_VECTOR_GET_PTR(buffers[0]);
+	int factor = *(int*)args;
+	unsigned int i;
+	for (i = 0; i < n; i++)
+	{
+		if (val[i] != i*factor)
+		{
+			vector_config.copy_failed = 1;
+			return;
+		}
+		val[i] = -val[i];
+	}
+}
+
+int
+main(void)
+{
+	data_interface_test_summary *summary;
+	struct starpu_conf conf =
+	{
+		.ncpus   = -1,
+		.ncuda   = 2,
+		.nopencl = 1
+	};
+
+	if (starpu_init(&conf) == -ENODEV)
+		goto enodev;
+
+	register_data();
+
+	summary = run_tests(&vector_config);
+	if (!summary)
+		exit(EXIT_FAILURE);
+
+	unregister_data();
+
+	starpu_shutdown();
+
+	data_interface_test_summary_print(stderr, summary);
+
+	return data_interface_test_summary_success(summary);
+
+enodev:
+	return STARPU_TEST_SKIPPED;
+}

tests/datawizard/interfaces/vector/test_vector_opencl.c

@@ -0,0 +1,130 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+#include <config.h>
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include "../test_interfaces.h"
+
+#define KERNEL_LOCATION "tests/datawizard/interfaces/vector/test_vector_opencl_kernel.cl"
+extern struct test_config vector_config;
+static struct starpu_opencl_program opencl_program;
+
+void
+test_vector_opencl_func(void *buffers[], void *args)
+{
+	STARPU_SKIP_IF_VALGRIND;
+
+	int id, devid, ret;
+	int factor = *(int *) args;
+
+        cl_int             err;
+	cl_kernel          kernel;
+	cl_command_queue   queue;
+	cl_event           event;
+
+	ret = starpu_opencl_load_opencl_from_file(KERNEL_LOCATION, &opencl_program, NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+
+	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	cl_mem val = (cl_mem)STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
+
+	cl_context context;
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+	starpu_opencl_get_context(devid, &context);
+
+	cl_mem fail = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+		sizeof(int), &vector_config.copy_failed, &err);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_program,
+					"test_vector_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	int nargs;
+	nargs = starpu_opencl_set_kernel_args(&err, &kernel,
+					      sizeof(val), &val,
+					      sizeof(n), &n,
+					      sizeof(fail), &fail,
+					      sizeof(factor), &factor,
+					      0);
+
+	if (nargs != 4)
+	{
+		fprintf(stderr, "Failed to set argument #%d\n", err);
+		STARPU_OPENCL_REPORT_ERROR(err);
+	}
+			
+	{
+		size_t global=n;
+		size_t local;
+                size_t s;
+                cl_device_id device;
+
+                starpu_opencl_get_device(devid, &device);
+
+                err = clGetKernelWorkGroupInfo (kernel,
+						device,
+						CL_KERNEL_WORK_GROUP_SIZE,
+						sizeof(local),
+						&local,
+						&s);
+                if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+
+                if (local > global)
+			local = global;
+
+		err = clEnqueueNDRangeKernel(queue,
+					kernel,
+					1,
+					NULL,
+					&global,
+					&local,
+					0,
+					NULL,
+					&event);
+
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	err = clEnqueueReadBuffer(queue,
+				  fail,
+				  CL_TRUE,
+				  0, 
+				  sizeof(int),
+				  &vector_config.copy_failed,
+				  0,
+				  NULL,
+				  NULL);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+        starpu_opencl_unload_opencl(&opencl_program);
+}

tests/datawizard/interfaces/vector/test_vector_opencl_kernel.cl

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+__kernel void test_vector_opencl(__global int *val,
+				 unsigned int nx,
+				 __global int *err,
+				 int factor)
+{
+        const int i = get_global_id(0);
+        if (i >=  nx)
+		return;
+
+	if (val[i] != i * factor)
+		*err = 1;
+	else
+		val[i] = - val[i];
+}