Add standard test for the matrix interface

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/matrix/matrix_opencl_kernel.cl \
 	datawizard/interfaces/variable/variable_opencl_kernel.cl \
 	datawizard/interfaces/vector/test_vector_opencl_kernel.cl \
 	datawizard/interfaces/multiformat/multiformat_types.h \
@@ -158,6 +159,7 @@ noinst_PROGRAMS =				\
 	datawizard/increment_redux_v2		\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
+	datawizard/interfaces/matrix/matrix_interface \
 	datawizard/interfaces/multiformat/multiformat_interface \
 	datawizard/interfaces/variable/variable_interface    \
 	datawizard/interfaces/vector/test_vector_interface   \
@@ -286,6 +288,25 @@ nobase_STARPU_OPENCL_DATA_DATA += \
 	datawizard/interfaces/vector/test_vector_opencl_kernel.cl
 endif
 
+####################
+# Matrix interface #
+####################
+datawizard_interfaces_matrix_matrix_interface_SOURCES= \
+	datawizard/interfaces/test_interfaces.c        \
+	datawizard/interfaces/matrix/matrix_interface.c
+
+if STARPU_USE_CUDA
+datawizard_interfaces_matrix_matrix_interface_SOURCES+= \
+	datawizard/interfaces/matrix/matrix_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_matrix_matrix_interface_SOURCES+= \
+	datawizard/interfaces/matrix/matrix_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA+= \
+	datawizard/interfaces/matrix/matrix_opencl_kernel.cl
+endif
+
 datawizard_interfaces_multiformat_multiformat_interface_SOURCES =           \
 	datawizard/interfaces/test_interfaces.c                             \
 	datawizard/interfaces/multiformat/multiformat_interface.c           \
@@ -321,4 +342,3 @@ datawizard_interfaces_variable_variable_interface_SOURCES+= \
 nobase_STARPU_OPENCL_DATA_DATA += \
 	datawizard/interfaces/variable/variable_opencl_kernel.cl
 endif
-

tests/datawizard/interfaces/matrix/matrix_cuda.cu

@@ -0,0 +1,71 @@
+/* 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 matrix_config;
+
+__global__ void matrix_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_matrix_cuda_func(void *buffers[], void *args)
+{
+	int factor;
+	int *ret;
+	int *val;
+	cudaError_t error;
+	unsigned int nx, ny, n;
+
+	nx = STARPU_MATRIX_GET_NX(buffers[0]);
+	ny = STARPU_MATRIX_GET_NY(buffers[0]);
+	n = nx * ny;
+	unsigned threads_per_block = 64;
+	unsigned nblocks = (n + threads_per_block-1) / threads_per_block;
+	factor = *(int *) args;
+	val = (int *) STARPU_MATRIX_GET_PTR(buffers[0]);
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	error = cudaMemcpy(ret,
+			   &matrix_config.copy_failed,
+			   sizeof(int),
+			   cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+        matrix_cuda<<<nblocks,threads_per_block,2,starpu_cuda_get_local_stream()>>>(val, n, ret, factor);
+
+	error = cudaMemcpy(&matrix_config.copy_failed,
+			   ret,
+			   sizeof(int),
+			   cudaMemcpyDeviceToHost);
+	
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/matrix/matrix_interface.c

@@ -0,0 +1,140 @@
+/* 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"
+#include "../../../common/helper.h"
+
+#define WIDTH  16
+#define HEIGHT 16
+
+#ifdef STARPU_USE_CPU
+static void test_matrix_cpu_func(void *buffers[], void *args);
+#endif /* !STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+extern void test_matrix_cuda_func(void *buffers[], void *_args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_matrix_opencl_func(void *buffers[], void *args);
+#endif
+
+
+static starpu_data_handle_t matrix_handle;
+static starpu_data_handle_t matrix2_handle;
+
+struct test_config matrix_config = {
+#ifdef STARPU_USE_CPU
+	.cpu_func      = test_matrix_cpu_func,
+#endif /* ! STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+	.cuda_func     = test_matrix_cuda_func,
+#endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_func   = test_matrix_opencl_func,
+#endif
+	.handle        = &matrix_handle,
+	.dummy_handle  = &matrix2_handle,
+	.copy_failed   = 0,
+	.name          = "matrix_interface"
+};
+
+static int matrix[WIDTH * HEIGHT];
+static int matrix2[WIDTH * HEIGHT];
+
+static void
+register_data(void)
+{
+	int i;
+	int size = WIDTH * HEIGHT;
+	for (i = 0; i < size; i++)
+		matrix[i] = i;
+
+	starpu_matrix_data_register(&matrix_handle,
+				    0,
+				    (uintptr_t) matrix,
+				    WIDTH, /* ld */
+				    WIDTH,
+				    HEIGHT,
+				    sizeof(matrix[0]));
+	starpu_matrix_data_register(&matrix2_handle,
+				    0,
+				    (uintptr_t) matrix2,
+				    WIDTH, /* ld */
+				    WIDTH,
+				    HEIGHT,
+				    sizeof(matrix[0]));
+}
+
+static void
+unregister_data(void)
+{
+	starpu_data_unregister(matrix_handle);
+	starpu_data_unregister(matrix2_handle);
+}
+
+static void
+test_matrix_cpu_func(void *buffers[], void *args)
+{
+	int *val;
+	int factor;
+	int i;
+	unsigned int nx, ny;
+
+	nx = STARPU_MATRIX_GET_NX(buffers[0]);
+	ny = STARPU_MATRIX_GET_NY(buffers[0]);
+	val = (int *) STARPU_MATRIX_GET_PTR(buffers[0]);
+	factor = *(int *) args;
+
+	for (i = 0; i < nx*ny; i++)
+	{
+		if (val[i] != i * factor)
+		{
+			matrix_config.copy_failed = 1;
+			return;
+		}
+		val[i] *= -1;
+	}
+}
+
+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(&matrix_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/matrix/matrix_opencl.c

@@ -0,0 +1,125 @@
+/* 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/matrix/matrix_opencl_kernel.cl"
+
+extern struct test_config matrix_config;
+static struct starpu_opencl_program matrix_program;
+
+void test_matrix_opencl_func(void *buffers[], void *args)
+{
+	int id, devid, factor;
+	unsigned int n;
+
+        cl_int             err;
+	cl_kernel          kernel;
+	cl_command_queue   queue;
+	cl_event           event;
+	cl_context         context;
+	cl_mem             val, fail;
+
+	starpu_opencl_load_opencl_from_file(KERNEL_LOCATION,
+					    &matrix_program,
+					    NULL);
+
+	factor = *(int *)args;
+	n = STARPU_MATRIX_GET_NX(buffers[0]);
+	n*= STARPU_MATRIX_GET_NY(buffers[0]);
+	val = (cl_mem)STARPU_MATRIX_GET_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,
+					&matrix_program,
+					"matrix_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	fail = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+		sizeof(int), &matrix_config.copy_failed, &err);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	/* Setting args */
+	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)
+		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),
+				  &matrix_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(&matrix_program);
+}
+

tests/datawizard/interfaces/matrix/matrix_opencl_kernel.cl

@@ -0,0 +1,31 @@
+/* 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 matrix_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 = i;
+	else
+		val[i] *= -1;
+}
+