Variable interface : standard test.

tests/Makefile.am

@@ -27,6 +27,7 @@ EXTRA_DIST =					\
 	datawizard/sync_and_notify_data_opencl_codelet.cl\
 	coverage/coverage.sh			\
 	datawizard/interfaces/test_interfaces.h	\
+	datawizard/interfaces/variable/variable_opencl_kernel.cl \
 	datawizard/interfaces/vector/test_vector_opencl_kernel.cl \
 	datawizard/interfaces/multiformat/multiformat_types.h
 
@@ -156,6 +157,7 @@ noinst_PROGRAMS =				\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
 	datawizard/interfaces/multiformat/multiformat_interface \
+	datawizard/interfaces/variable/variable_interface    \
 	datawizard/interfaces/vector/test_vector_interface   \
 	errorcheck/starpu_init_noworker		\
 	errorcheck/invalid_blocking_calls	\
@@ -301,3 +303,20 @@ nobase_STARPU_OPENCL_DATA_DATA +=
 	datawizard/interfaces/multiformat/multiformat_opencl_kernel.cl                     \
 	datawizard/interfaces/multiformat/multiformat_conversion_codelets_opencl_kernel.cl
 endif
+
+datawizard_interfaces_variable_variable_interface_SOURCES=   \
+	datawizard/interfaces/test_interfaces.c              \
+	datawizard/interfaces/variable/variable_interface.c
+
+if STARPU_USE_CUDA
+datawizard_interfaces_variable_variable_interface_SOURCES+= \
+	datawizard/interfaces/variable/variable_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_variable_variable_interface_SOURCES+= \
+	datawizard/interfaces/variable/variable_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	datawizard/interfaces/variable/variable_opencl_kernel.cl
+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,106 @@
+/* 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"
+
+static int variable;
+static starpu_data_handle_t variable_handle;
+
+/* Codelets */
+#ifdef STARPU_USE_CPU
+static void test_variable_cpu_func(void *buffers[], void *args);
+#endif
+#ifdef STARPU_USE_CUDA
+static void test_variable_cuda_func(void *buffers[], void *args);
+#endif
+#ifdef STARPU_USE_OPENCL
+static 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,
+	.copy_failed = 0,
+	.name        = "variable_interface"
+};
+
+static void
+test_variable_cpu_func(void *buffers[], void *args)
+{
+	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;
+	starpu_variable_data_register(&variable_handle, 0,
+				      (uintptr_t) &variable, sizeof(variable));
+}
+
+static
+void unregister_data(void)
+{
+	starpu_data_unregister(variable_handle);
+}
+
+int
+main(void)
+{
+	data_interface_test_summary *summary;
+
+	struct starpu_conf conf =
+	{
+		.ncpus = -1,
+		.ncuda = 2,
+		.nopencl = 1
+	};
+
+	starpu_init(&conf);
+
+	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,124 @@
+/* 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/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)
+{
+	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);
+
+	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;
+}
+