CSR interface : add standard test

tests/Makefile.am

@@ -161,6 +161,7 @@ noinst_PROGRAMS =				\
 	datawizard/increment_redux_v2		\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
+	datawizard/interfaces/csr/csr_interface \
 	datawizard/interfaces/matrix/matrix_interface \
 	datawizard/interfaces/multiformat/multiformat_interface \
 	datawizard/interfaces/multiformat/advanced/multiformat_cuda_opencl \
@@ -272,9 +273,26 @@ core_multiformat_handle_conversion_SOURCES = \
 core_multiformat_data_release_SOURCES = \
 	core/multiformat_data_release.c
 
-##############
-# Interfaces #
-##############
+#################
+# CSR interface #
+#################
+datawizard_interfaces_csr_csr_interface_SOURCES= \
+	datawizard/interfaces/test_interfaces.c  \
+	datawizard/interfaces/csr/csr_interface.c
+
+if STARPU_USE_CUDA
+datawizard_interfaces_csr_csr_interface_SOURCES+= \
+	datawizard/interfaces/csr/csr_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_csr_csr_interface_SOURCES+= \
+	datawizard/interfaces/csr/csr_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	datawizard/interfaces/csr/csr_opencl_kernel.cl
+endif
+
+
 datawizard_interfaces_vector_test_vector_interface_SOURCES =               \
 	datawizard/interfaces/vector/test_vector_interface.c               \
 	datawizard/interfaces/test_interfaces.c

tests/datawizard/interfaces/csr/csr_cuda.cu

@@ -0,0 +1,68 @@
+/* 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 csr_config;
+
+__global__ void csr_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+1)*factor)
+		*err = 1;
+	else
+		nzval[i] = -nzval[i];
+}
+
+extern "C" void test_csr_cuda_func(void *buffers[], void *args)
+{
+	int factor;
+	int *ret;
+	int *val;
+	cudaError_t error;
+	uint32_t nnz = STARPU_CSR_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_CSR_GET_NZVAL(buffers[0]);
+
+	error = cudaMalloc(&ret, sizeof(int));
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	error = cudaMemcpy(ret,
+			   &csr_config.copy_failed,
+			   sizeof(int),
+			   cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+        csr_cuda<<<nblocks,threads_per_block,2,starpu_cuda_get_local_stream()>>> (val, nnz, ret, factor);
+
+	error = cudaMemcpy(&csr_config.copy_failed,
+			   ret,
+			   sizeof(int),
+			   cudaMemcpyDeviceToHost);
+	
+	cudaFree(ret);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}

tests/datawizard/interfaces/csr/csr_interface.c

@@ -0,0 +1,167 @@
+/* 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  8
+#define HEIGHT 4
+#define SIZE   (WIDTH * HEIGHT)
+#define NNZ    (SIZE-1)
+
+#ifdef STARPU_USE_CPU
+static void test_csr_cpu_func(void *buffers[], void *args);
+#endif /* !STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+extern void test_csr_cuda_func(void *buffers[], void *_args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_csr_opencl_func(void *buffers[], void *args);
+#endif
+
+
+static int nzval[NNZ];
+static int nzval2[NNZ];
+
+static uint32_t colind[NNZ];
+static uint32_t colind2[NNZ];
+
+static uint32_t rowptr[HEIGHT+1];
+static uint32_t rowptr2[HEIGHT+1];
+
+static starpu_data_handle_t csr_handle;
+static starpu_data_handle_t csr2_handle;
+
+struct test_config csr_config =
+{
+#ifdef STARPU_USE_CPU
+	.cpu_func      = test_csr_cpu_func,
+#endif /* ! STARPU_USE_CPU */
+#ifdef STARPU_USE_CUDA
+	.cuda_func     = test_csr_cuda_func,
+#endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_func   = test_csr_opencl_func,
+#endif
+	.handle        = &csr_handle,
+	.dummy_handle  = &csr2_handle,
+	.copy_failed   = 0,
+	.name          = "csr_interface"
+};
+
+static void
+register_data(void)
+{
+	int i;
+	for (i = 1; i < SIZE; i++)
+	{
+		nzval[i-1] = i;
+		nzval2[i-1] = 42;
+
+		colind[i-1] = i % WIDTH;
+		colind2[i-1] = colind[i];
+	}
+
+	rowptr[0] = 1;
+	rowptr2[0] = 1;
+	for (i = 1; i < HEIGHT; i++)
+	{
+		rowptr[i] = i * WIDTH;
+		rowptr2[i] = rowptr[i];
+	}
+	rowptr[HEIGHT] = NNZ + 1;
+	rowptr2[HEIGHT] = rowptr[HEIGHT];
+
+	starpu_csr_data_register(&csr_handle,
+				 0,
+				 NNZ,
+				 HEIGHT,
+				 (uintptr_t) nzval,
+				 colind,
+				 rowptr,
+				 0,
+				 sizeof(nzval[0]));
+	starpu_csr_data_register(&csr2_handle,
+				 0,
+				 NNZ,
+				 HEIGHT,
+				 (uintptr_t) nzval2,
+				 colind2,
+				 rowptr2,
+				 0,
+				 sizeof(nzval2[0]));
+}
+
+static void
+unregister_data(void)
+{
+	starpu_data_unregister(csr_handle);
+	starpu_data_unregister(csr2_handle);
+}
+
+static void
+test_csr_cpu_func(void *buffers[], void *args)
+{
+	int *val;
+	int factor;
+	int i;
+
+	uint32_t nnz = STARPU_CSR_GET_NNZ(buffers[0]);
+	val = (int *) STARPU_CSR_GET_NZVAL(buffers[0]);
+	factor = *(int *) args;
+
+	for (i = 0; i < nnz; i++)
+	{
+		if (val[i] != (i+1) * factor)
+		{
+			csr_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(&csr_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/csr/csr_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/csr/csr_opencl_kernel.cl"
+extern struct test_config csr_config;
+static struct starpu_opencl_program opencl_program;
+
+void
+test_csr_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_CSR_GET_NNZ(buffers[0]);
+	cl_mem nzval = (cl_mem)STARPU_CSR_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), &csr_config.copy_failed, &err);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_program,
+					"test_csr_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),
+				  &csr_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/csr/csr_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_csr_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+1) * factor)
+		*err = 1;
+	else
+		val[i] = - val[i];
+}