Generic test for the interfaces.

tests/Makefile.am

@@ -147,6 +147,7 @@ noinst_PROGRAMS =				\
 	datawizard/increment_redux_v2		\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
+	datawizard/interfaces/test_interfaces   \
 	errorcheck/starpu_init_noworker		\
 	errorcheck/invalid_blocking_calls	\
 	errorcheck/invalid_tasks		\
@@ -166,7 +167,7 @@ noinst_PROGRAMS =				\
 	parallel_tasks/parallel_kernels		\
 	parallel_tasks/parallel_kernels_spmd	\
 	perfmodels/regression_based		\
-	perfmodels/non_linear_regression_based
+	perfmodels/non_linear_regression_based 
 
 if STARPU_HAVE_WINDOWS
 check_PROGRAMS = $(noinst_PROGRAMS)
@@ -246,3 +247,19 @@ BUILT_SOURCES += 						\
 	datawizard/sync_and_notify_data_gordon_kernels.spuelf	\
 	microbenchs/null_kernel_gordon.spuelf
 endif
+
+datawizard_interfaces_test_interfaces_SOURCES =               \
+	datawizard/interfaces/test_interfaces.c     \
+	datawizard/interfaces/test_vector_interface.c
+
+if STARPU_USE_CUDA
+datawizard_interfaces_test_interfaces_SOURCES +=                \
+	datawizard/interfaces/test_vector_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_test_interfaces_SOURCES +=                           \
+	datawizard/interfaces/test_vector_opencl.c 
+nobase_STARPU_OPENCL_DATA_DATA += \
+	datawizard/interfaces/test_vector_opencl_kernel.cl
+endif

tests/datawizard/interfaces/test_interfaces.c

@@ -0,0 +1,390 @@
+/* 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>
+#ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
+#endif
+
+#include "test_interfaces.h"
+
+/* TODO :
+- OpenCL to OpenCL support
+- RAM to RAM ?
+- Asynchronous vs synchronous
+- Better error messages
+*/
+
+/* Interfaces to test */
+extern struct test_config vector_config;
+
+static struct test_config *tests[] = {
+	&vector_config,
+	NULL
+};
+
+static struct test_config *current_config;
+
+/*
+ * This variable has to be either -1 or 1.
+ * The kernels should check that the ith value stored in the data interface is
+ * equal to i, if factor == 1, or -i, if factor == -1.
+ */
+static int factor = -1;
+
+/*
+ * Creates a complete task, only knowing on what device it should be executed.
+ * Note that the global variable <current_config> is heavily used here.
+ * Arguments :
+ *	- taskp : a pointer to a valid task
+ *	- type : STARPU_{CPU,CUDA,OPENCL}_WORKER. Gordon is not supported.
+ *      - id   : -1 if you dont care about the device where the task will be 
+ *		 executed, as long as it has the right type.
+ *		 >= 0 if you want to make sure the task will be executed on the
+ *		 idth device that has the specified type.
+ * Return values :
+ * 	-ENODEV
+ *	0 : success.
+ */
+static int
+create_task(struct starpu_task **taskp, enum starpu_archtype type, int id)
+{
+	static int cpu_workers[STARPU_MAXCPUS];
+	static int cuda_workers[STARPU_MAXCUDADEVS];
+	static int opencl_workers[STARPU_MAXOPENCLDEVS];
+
+	static int n_cpus = -1;
+	static int n_cudas = -1;
+	static int n_opencls = -1;
+
+	if (n_cpus == -1) /* First time here */
+	{
+		/* XXX Dont check them all at once. */
+		/* XXX Error checking */
+		n_cpus = starpu_worker_get_ids_by_type(STARPU_CPU_WORKER,
+							cpu_workers,
+							STARPU_MAXCPUS);
+		n_cudas = starpu_worker_get_ids_by_type(STARPU_CUDA_WORKER,
+							cuda_workers,
+							STARPU_MAXCUDADEVS);
+		n_opencls = starpu_worker_get_ids_by_type(STARPU_OPENCL_WORKER,
+							opencl_workers,
+							STARPU_MAXOPENCLDEVS);
+	}
+
+	int workerid;
+	static struct starpu_codelet_t cl;
+	cl.nbuffers = 1;
+
+	switch (type)
+	{
+		case STARPU_CPU_WORKER:
+			if (id != -1)
+			{
+				if (id >= n_cpus)
+				{
+					fprintf(stderr, "Not enough CPU workers\n");
+					return -ENODEV;
+				}
+				workerid = *(cpu_workers + id);
+			}
+			cl.where = STARPU_CPU;
+			cl.cpu_func = current_config->cpu_func;
+			break;
+#ifdef STARPU_USE_CUDA
+		case STARPU_CUDA_WORKER:
+			if (id != -1)
+			{
+				if (id >= n_cudas)
+				{
+					fprintf(stderr, "Not enough CUDA workers\n");
+					return -ENODEV;
+				}
+				workerid = cuda_workers[id];
+			}
+			cl.where = STARPU_CUDA;
+			cl.cuda_func = current_config->cuda_func;
+			break;
+#endif /* !STARPU_USE_CUDA */
+#ifdef STARPU_USE_OPENCL
+		case STARPU_OPENCL_WORKER:
+			if (id != -1)
+			{
+				if (id >= n_opencls)
+				{
+					fprintf(stderr, "Not enough OpenCL workers\n");
+					return -ENODEV;
+				}
+				workerid = *(opencl_workers + id);
+			}
+			cl.where = STARPU_OPENCL;
+			cl.opencl_func = current_config->opencl_func;
+			break;
+#endif /* ! STARPU_USE_OPENCL */
+		default:
+			return -ENODEV;
+	}
+
+
+	struct starpu_task *task = starpu_task_create();
+	task->synchronous = 1;
+	task->cl = &cl;
+	task->buffers[0].handle = *(current_config->register_func());
+	task->buffers[0].mode = STARPU_RW;
+	if (id != -1)
+	{
+		task->execute_on_a_specific_worker = 1;
+		task->workerid = workerid;
+	}
+	factor = -factor;
+	task->cl_arg = &factor;
+	task->cl_arg_size = sizeof(&factor);
+
+	*taskp = task;
+	return 0;
+}
+
+/*
+ * <device1>_to_<device2> functions.
+ * They all create and submit a task that has to be executed on <device2>,
+ * forcing a copy between <device1> and <device2>.
+ * XXX : could we sometimes use starp_insert_task() ? It seems hars because we
+ * need to set the execute_on_a_specific_worker field...
+ */
+#ifdef STARPU_USE_CUDA
+static int
+ram_to_cuda(void)
+{
+	int err;
+	struct starpu_task *task;
+
+	err = create_task(&task, STARPU_CUDA_WORKER, 0);
+	if (err != 0)
+	{
+		fprintf(stderr, "Could not create task\n");
+		return 1;
+	}
+
+	err = starpu_task_submit(task);
+	if (err != 0)
+	{
+		fprintf(stderr, "Fail : %s\n", strerror(-err));
+		return 1;
+	}
+
+	fprintf(stderr, "[%s] : %d\n", __func__, current_config->copy_failed);
+	return current_config->copy_failed;
+}
+
+static int
+cuda_to_cuda(void)
+{
+	int err;
+	struct starpu_task *task;
+
+	err = create_task(&task, STARPU_CUDA_WORKER, 1);
+	if (err != 0)
+	{
+		return 1;
+	}
+
+	err = starpu_task_submit(task);
+	if (err != 0)
+	{
+		return 1;
+	}
+
+	fprintf(stderr, "[%s] : %d\n", __func__, current_config->copy_failed);
+	return current_config->copy_failed;
+}
+
+static int
+cuda_to_ram(void)
+{
+	int err;
+	struct starpu_task *task;
+
+	err = create_task(&task, STARPU_CPU_WORKER, -1);
+	if (err != 0)
+	{
+		fprintf(stderr, "Could not create the task\n");	
+		return 1;
+	}
+
+	err = starpu_task_submit(task);
+	if (err != 0)
+	{
+		fprintf(stderr, "Fail : %s\n", strerror(-err));
+		return 1;
+	}
+
+	fprintf(stderr, "[%s] : %d\n", __func__, current_config->copy_failed);
+	return current_config->copy_failed;
+}
+#endif /* !STARPU_USE_CUDA */
+
+#ifdef STARPU_USE_OPENCL
+static int
+ram_to_opencl()
+{
+	int err;
+	struct starpu_task *task;
+
+	err = create_task(&task, STARPU_OPENCL_WORKER, 0);
+	if (err != 0)
+	{
+		fprintf(stderr, "Could not create the task\n");
+		return 1;
+	}
+
+	err = starpu_task_submit(task);
+	if (err != 0)
+	{
+		fprintf(stderr, "Fail : %s\n", strerror(-err));
+		return 1;
+	}
+
+	fprintf(stderr, "[%s] : %d\n", __func__, current_config->copy_failed);
+	return current_config->copy_failed;
+}
+
+static int
+opencl_to_ram()
+{
+	int err;
+	struct starpu_task *task;
+
+	err = create_task(&task, STARPU_CPU_WORKER, -1);
+	if (err != 0)
+	{
+		fprintf(stderr, "Could not create the task\n");
+		return 1;
+	}
+
+	err = starpu_task_submit(task);
+	if (err != 0)
+	{
+		fprintf(stderr, "Fail : %s\n", strerror(-err));
+		return 1;
+	}
+
+	fprintf(stderr, "[%s] : %d\n", __func__, current_config->copy_failed);
+	return current_config->copy_failed;
+}
+#endif /* !STARPU_USE_OPENCL */
+/* End of the <device1>_to_<device2> functions. */
+
+static int
+run(void)
+{
+	int err;
+#ifdef STARPU_USE_CUDA
+	/* RAM -> CUDA -> CUDA -> RAM */
+	err = ram_to_cuda();
+	if (err != 0)
+	{
+		fprintf(stderr, "RAM to CUDA failed\n");
+		return 1;
+	}
+
+#ifdef HAVE_CUDA_MEMCPY_PEER
+	err = cuda_to_cuda();
+	if (err != 0)
+	{
+		fprintf(stderr, "CUDA to RAM failed\n");
+		return 1;
+	}
+#endif /* !HAVE_CUDA_MEMCPY_PEER */
+
+	err = cuda_to_ram();
+	if (err != 0)
+	{
+		fprintf(stderr, "CUDA to RAM failed\n");
+		return 1;
+	}
+#endif /* !STARPU_USE_CUDA */
+
+#if STARPU_USE_OPENCL
+	/* RAM -> OpenCL -> RAM */
+	err = ram_to_opencl();
+	if (err != 0)
+	{
+		fprintf(stderr, "RAM to OpenCL failed\n");
+		return 1;
+	}
+
+	err = opencl_to_ram();
+	if (err != 0)
+	{
+		fprintf(stderr, "OpenCL to RAM failed\n");
+		return 1;
+	}
+#endif /* !STARPU_USE_OPENCL */
+
+	return 0;
+}
+
+static int
+load_conf(struct test_config *config)
+{
+	if (!config ||
+	    !config->cpu_func ||
+#ifdef STARPU_USE_CUDA
+	    !config->cuda_func ||
+#endif
+#ifdef STARPU_USE_OPENCL
+	    !config->opencl_func ||
+#endif
+	    !config->register_func)
+	{
+		return 1;
+	}
+	
+	current_config = config;
+	return 0;
+}
+
+int
+main(void)
+{
+	int i;
+	int err;
+
+	err = starpu_init(NULL);
+	if (err != 0)
+	{
+		fprintf(stderr, "starpu_init failed, not running the tests\n");
+		return EXIT_FAILURE;
+	}
+
+	for (i = 0; tests[i] != NULL; ++i)
+	{
+		err = load_conf(tests[i]);
+		if (err != 0)
+		{
+			fprintf(stderr, "Skipping test, invalid conf\n");
+			continue;
+		}
+
+		err = run();
+		if (err != 0)
+			fprintf(stderr, "%s : FAIL\n", current_config->name);
+		else
+			fprintf(stderr, "%s : OK\n", current_config->name);
+	}
+	
+	starpu_shutdown();
+	return EXIT_SUCCESS;
+}

tests/datawizard/interfaces/test_interfaces.h

@@ -0,0 +1,45 @@
+/* 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_INTERFACES_H
+#define TEST_INTERFACES_H 
+
+typedef starpu_data_handle* (*register_func_t) (void);
+
+struct test_config {
+	/* Returns a valid handle to a piece of data registered by StarPU */
+	register_func_t register_func;
+
+	/* StarPU codelets. The following functions should :
+	 * 1) Check that the values are correct
+	 * 2) Negate every element
+	 */
+	starpu_cpu_func_t cpu_func;
+#ifdef STARPU_USE_CUDA
+	starpu_cuda_func_t cuda_func;
+#endif
+#ifdef STARPU_USE_OPENCL
+	starpu_opencl_func_t opencl_func;
+#endif
+
+	/* The previous codelets must update this field at the end of their
+	 * execution. copy_failed must be 1 if the copy failed, 0 otherwise. */
+	int copy_failed;
+
+	/* A human-readable name for the test */
+	const char *name;
+};
+
+#endif /* !TEST_INTERFACES_H */

tests/datawizard/interfaces/test_vector_cuda.cu

@@ -0,0 +1,65 @@
+/* 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/test_vector_interface.c

@@ -0,0 +1,88 @@
+/* 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 "test_interfaces.h"
+
+/* Prototypes */
+static starpu_data_handle *register_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 *vector_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
+	.register_func = register_data,
+	.copy_failed   = 0,
+	.name          = "vector_interface"
+};
+
+int n = 16;
+int *vector;
+
+static starpu_data_handle*
+register_data(void)
+{
+	if (vector_handle)
+		return vector_handle;
+
+	/* Initializing data */
+	int i;
+	vector = malloc(n * sizeof(*vector));
+	if (!vector)
+		return NULL;
+	for (i = 0; i < n; i++)
+		vector[i] = i;
+
+	/* Registering data */
+	vector_handle = malloc(sizeof(*vector_handle));
+	if (!vector_handle)
+		return NULL;
+	starpu_vector_data_register(vector_handle,
+                                    0,
+                                    (uintptr_t)vector,
+                                     n,
+                                     sizeof(int));
+	return vector_handle;
+}
+
+static void test_vector_cpu_func(void *buffers[], void *args)
+{
+	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) {
+			fprintf(stderr, "HI %d => %d\n", i, val[i]);
+			vector_config.copy_failed = 1;
+			return;
+		}
+		val[i] = -val[i];
+	}
+}

tests/datawizard/interfaces/test_vector_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 <starpu.h>
+#include <starpu_opencl.h>
+#include "test_interfaces.h"
+
+#define KERNEL_LOCATION "tests/datawizard/interfaces/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)
+{
+	int id, devid;
+	int factor = *(int *) args;
+
+        cl_int             err;
+	cl_kernel          kernel;
+	cl_command_queue   queue;
+	cl_event           event;
+
+	starpu_opencl_load_opencl_from_file(KERNEL_LOCATION, &opencl_program, NULL);
+
+	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	cl_mem val = (cl_mem)STARPU_VECTOR_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), &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);
+
+	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),
+				  &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/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];
+}