BCSR interface : add standard test

include/starpu_data_interfaces.h

@@ -327,6 +327,10 @@ void starpu_bcsr_data_register(starpu_data_handle_t *handle, uint32_t home_node,
 		uintptr_t nzval, uint32_t *colind, uint32_t *rowptr, uint32_t firstentry, uint32_t r, uint32_t c, size_t elemsize);
 
 
+#define STARPU_BCSR_GET_NNZ(interface)        (((struct starpu_bcsr_interface *)(interface))->nnz)
+#define STARPU_BCSR_GET_NZVAL(interface)      (((struct starpu_bcsr_interface *)(interface))->nzval)
+#define STARPU_BCSR_GET_COLIND(interface)     (((struct starpu_bcsr_interface *)(interface))->colind)
+#define STARPU_BCSR_GET_ROWPTR(interface)     (((struct starpu_bcsr_interface *)(interface))->rowptr)
 uint32_t starpu_bcsr_get_nnz(starpu_data_handle_t handle);
 uint32_t starpu_bcsr_get_nrow(starpu_data_handle_t handle);
 uint32_t starpu_bcsr_get_firstentry(starpu_data_handle_t handle);

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/bcsr/bcsr_opencl_kernel.cl \
 	datawizard/interfaces/matrix/matrix_opencl_kernel.cl \
 	datawizard/interfaces/variable/variable_opencl_kernel.cl \
 	datawizard/interfaces/vector/test_vector_opencl_kernel.cl \
@@ -166,6 +167,7 @@ noinst_PROGRAMS =				\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
 	datawizard/interfaces/block/block_interface \
+	datawizard/interfaces/bcsr/bcsr_interface \
 	datawizard/interfaces/csr/csr_interface \
 	datawizard/interfaces/matrix/matrix_interface \
 	datawizard/interfaces/multiformat/multiformat_interface \
@@ -298,6 +300,25 @@ nobase_STARPU_OPENCL_DATA_DATA += \
 	datawizard/interfaces/block/block_opencl_kernel.cl
 endif
 
+##################
+# BSCR interface #
+##################
+datawizard_interfaces_bcsr_bcsr_interface_SOURCES= \
+	datawizard/interfaces/test_interfaces.c \
+	datawizard/interfaces/bcsr/bcsr_interface.c 
+
+if STARPU_USE_CUDA
+datawizard_interfaces_bcsr_bcsr_interface_SOURCES+= \
+	datawizard/interfaces/bcsr/bcsr_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_bcsr_bcsr_interface_SOURCES+= \
+	datawizard/interfaces/bcsr/bcsr_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	datawizard/interfaces/bcsr/bcsr_opencl_kernel.cl
+endif
+
 #################
 # CSR interface #
 #################

tests/datawizard/interfaces/bcsr/bcsr_cuda.cu

@@ -0,0 +1,70 @@
+/* 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 bcsr_config;
+
+__global__ void bcsr_cuda(int *nzval, uint32_t nnz, int *err, int factor)
+{
+        unsigned i =  blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (i >= nnz)
+		return;
+
+	if (nzval[i] != i*factor)
+		*err = 1;
+	else
+		nzval[i] = -nzval[i];
+}
+
+extern "C" void test_bcsr_cuda_func(void *buffers[], void *args)
+{
+	int factor;
+	int *ret;
+	int *val;
+	cudaError_t error;
+	uint32_t nnz = STARPU_BCSR_GET_NNZ(buffers[0]);
+	unsigned threads_per_block = 64;
+	unsigned nblocks = (nnz + threads_per_block-1) / threads_per_block;
+
+	factor = *(int *) args;
+	//val = (int *) starpu_bcsr_get_local_nzval((starpu_data_handle_t)buffers[0]);
+	val = (int *) STARPU_BCSR_GET_NZVAL(buffers[0]);
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	error = cudaMemcpy(ret,
+			   &bcsr_config.copy_failed,
+			   sizeof(int),
+			   cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+        bcsr_cuda<<<nblocks,threads_per_block,2,starpu_cuda_get_local_stream()>>>
+		(val, nnz, ret, factor);
+
+	error = cudaMemcpy(&bcsr_config.copy_failed,
+			   ret,
+			   sizeof(int),
+			   cudaMemcpyDeviceToHost);
+	
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/bcsr/bcsr_interface.c

@@ -0,0 +1,195 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ * 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 "../../../helper.h"
+
+/*
+ * XXX : These values should not be changed. If you really understand all that
+ * BCSR stuff, feel free to write a better example :)
+ */
+
+/* Size of the matrix */
+#define WIDTH          4
+#define HEIGHT         4
+#define SIZE           (WIDTH * HEIGHT)
+
+/* Size of the blocks */
+#define R              2
+#define C              2
+#define BLOCK_SIZE     (R*C)
+
+/* The matrix is simply 0 1 2... There are SIZE-1 non zero values... */
+#define NNZ            (SIZE-1)
+
+/* ... and SIZE/BLOCK_SIZE non zero blocks */
+#define NNZ_BLOCKS     (SIZE/BLOCK_SIZE)
+
+
+#ifdef STARPU_USE_CPU
+static void test_bcsr_cpu_func(void *buffers[], void *args);
+#endif /* !STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+extern void test_bcsr_cuda_func(void *buffers[], void *_args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_bcsr_opencl_func(void *buffers[], void *args);
+#endif
+
+
+static int nzval[NNZ];
+static int nzval2[NNZ];
+
+static uint32_t colind[NNZ_BLOCKS];
+static uint32_t colind2[NNZ_BLOCKS];
+
+static uint32_t rowptr[1+WIDTH/R];
+static uint32_t rowptr2[1+WIDTH/R];
+
+static starpu_data_handle_t bcsr_handle;
+static starpu_data_handle_t bcsr2_handle;
+
+
+struct test_config bcsr_config =
+{
+#ifdef STARPU_USE_CPU
+	.cpu_func      = test_bcsr_cpu_func,
+#endif /* !STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+	.cuda_func     = test_bcsr_cuda_func,
+#endif /* !STARPU_USE_CUDA */
+#ifdef STARPU_USE_OPENCL
+	.opencl_func   = test_bcsr_opencl_func,
+#endif /* !STARPU_USE_OPENCL */
+	.handle        = &bcsr_handle,
+	.dummy_handle  = &bcsr2_handle,
+	.copy_failed   = 0,
+	.name          = "bcsr_interface"
+};
+
+static void
+register_data(void)
+{
+	int i, j;
+
+	for (i = 0; i < SIZE; i++)
+		nzval[i] = i;
+
+	colind[0] = 0;
+	colind[1] = 2;
+	colind[2] = 0;
+	colind[3] = 2;
+
+	rowptr[0] = 0;
+	rowptr[1] = 2;
+	rowptr[2] = 4;
+	
+	starpu_bcsr_data_register(&bcsr_handle,
+				  0,
+				  NNZ_BLOCKS,
+				  HEIGHT/R,
+				  (uintptr_t) nzval,
+				  colind,
+				  rowptr,
+				  0,
+				  R,
+				  C,
+				  sizeof(nzval[0]));
+
+	starpu_bcsr_data_register(&bcsr2_handle,
+				  0,
+				  NNZ_BLOCKS,
+				  HEIGHT/R,
+				  (uintptr_t) nzval2,
+				  colind2,
+				  rowptr2,
+				  0,
+				  R,
+				  C,
+				  sizeof(nzval2[0]));
+}
+
+static void
+unregister_data(void)
+{
+	starpu_data_unregister(bcsr_handle);
+	starpu_data_unregister(bcsr2_handle);
+}
+
+static void
+test_bcsr_cpu_func(void *buffers[], void *args)
+{
+	int *val;
+	int factor;
+	int i;
+
+	uint32_t nnz = STARPU_BCSR_GET_NNZ(buffers[0]);
+	val = (int *) STARPU_BCSR_GET_NZVAL(buffers[0]);
+	factor = *(int *) args;
+
+	for (i = 0; i < nnz; i++)
+	{
+		if (val[i] != i * factor)
+		{
+			bcsr_config.copy_failed = 1;
+			return;
+		}
+		val[i] *= -1;
+	}
+
+	/* Check colind */
+	uint32_t *col = STARPU_BCSR_GET_COLIND(buffers[0]);
+	for (i = 0; i < NNZ_BLOCKS; i++)
+		if (col[i] != colind[i])
+			bcsr_config.copy_failed = 1;
+
+	/* Check rowptr */
+	uint32_t *row = STARPU_BCSR_GET_ROWPTR(buffers[0]);
+	for (i = 0; i < 1 + WIDTH/R; i++)
+		if (row[i] != rowptr[i])
+			bcsr_config.copy_failed = 1;
+}
+
+int
+main(void)
+{
+	data_interface_test_summary *summary;
+	struct starpu_conf conf =
+	{
+		.ncpus   = -1,
+		.ncuda   = 2,
+		.nopencl = 1
+	};
+
+	if (starpu_init(&conf) == -ENODEV)
+		return STARPU_TEST_SKIPPED;
+
+	register_data();
+
+	summary = run_tests(&bcsr_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/bcsr/bcsr_opencl.c

@@ -0,0 +1,126 @@
+/* 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/bcsr/bcsr_opencl_kernel.cl"
+extern struct test_config bcsr_config;
+static struct starpu_opencl_program opencl_program;
+
+void
+test_bcsr_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);
+
+	uint32_t nnz = STARPU_BCSR_GET_NNZ(buffers[0]);
+	cl_mem nzval = (cl_mem)STARPU_BCSR_GET_NZVAL(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), &bcsr_config.copy_failed, &err);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_program,
+					"test_bcsr_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	int nargs;
+	nargs = starpu_opencl_set_kernel_args(&err, &kernel,
+					      sizeof(nzval), &nzval,
+					      sizeof(nnz), &nnz,
+					      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 = nnz;
+		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),
+				  &bcsr_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/bcsr/bcsr_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_bcsr_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];
+}