Block interface : add standard test

tests/Makefile.am

@@ -163,6 +163,7 @@ noinst_PROGRAMS =				\
 	datawizard/increment_redux_v2		\
 	datawizard/handle_to_pointer		\
 	datawizard/lazy_allocation		\
+	datawizard/interfaces/block/block_interface \
 	datawizard/interfaces/csr/csr_interface \
 	datawizard/interfaces/matrix/matrix_interface \
 	datawizard/interfaces/multiformat/multiformat_interface \
@@ -275,6 +276,25 @@ core_multiformat_handle_conversion_SOURCES = \
 core_multiformat_data_release_SOURCES = \
 	core/multiformat_data_release.c
 
+###################
+# Block interface #
+###################
+datawizard_interfaces_block_block_interface_SOURCES= \
+	datawizard/interfaces/test_interfaces.c  \
+	datawizard/interfaces/block/block_interface.c
+
+if STARPU_USE_CUDA
+datawizard_interfaces_block_block_interface_SOURCES+= \
+	datawizard/interfaces/block/block_cuda.cu
+endif
+
+if STARPU_USE_OPENCL
+datawizard_interfaces_block_block_interface_SOURCES+= \
+	datawizard/interfaces/block/block_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	datawizard/interfaces/block/block_opencl_kernel.cl
+endif
+
 #################
 # CSR interface #
 #################

tests/datawizard/interfaces/block/block_cuda.cu

@@ -0,0 +1,80 @@
+/* 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 block_config;
+
+static __global__ void block_cuda(int *block,
+				  int nx, int ny, int nz,
+				  unsigned ldy, unsigned ldz,
+				  float factor, int *err)
+{
+        int i, j, k;
+	int val = 0;
+
+        for (k = 0; k < nz ;k++)
+	{
+                for (j = 0; j < ny ;j++)
+		{
+                        for(i = 0; i < nx ;i++)
+			{
+				if (block[(k*ldz)+(j*ldy)+i] != factor * val)
+				{
+					*err = 1;
+					return;
+				}
+				else
+				{
+					block[(k*ldz)+(j*ldy)+i] *= -1;
+					val++;
+				}
+			}
+                }
+        }
+}
+
+extern "C" void test_block_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, &block_config.copy_failed, sizeof(int), cudaMemcpyHostToDevice);
+	if (error != cudaSuccess)
+		STARPU_CUDA_REPORT_ERROR(error);
+
+	int nx = STARPU_BLOCK_GET_NX(buffers[0]);
+	int ny = STARPU_BLOCK_GET_NY(buffers[0]);
+	int nz = STARPU_BLOCK_GET_NZ(buffers[0]);
+        unsigned ldy = STARPU_BLOCK_GET_LDY(buffers[0]);
+        unsigned ldz = STARPU_BLOCK_GET_LDZ(buffers[0]);
+	int *block = (int *) STARPU_BLOCK_GET_PTR(buffers[0]);
+	int factor = *(int*) args;
+
+        block_cuda<<<1,1, 0, starpu_cuda_get_local_stream()>>>
+		(block, nx, ny, nz, ldy, ldz, factor, ret);
+	error = cudaMemcpy(&block_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/block/block_interface.c

@@ -0,0 +1,160 @@
+/* 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 "../../../helper.h"
+
+#define NX 16
+#define NY NX
+#define NZ NX
+
+/* Prototypes */
+static void register_data(void);
+static void unregister_data(void);
+static void test_block_cpu_func(void *buffers[], void *args);
+#ifdef STARPU_USE_CUDA
+extern void test_block_cuda_func(void *buffers[], void *_args);
+#endif
+#ifdef STARPU_USE_OPENCL
+extern void test_block_opencl_func(void *buffers[], void *args);
+#endif
+
+
+static starpu_data_handle_t block_handle;
+static starpu_data_handle_t block2_handle;
+
+struct test_config block_config =
+{
+	.cpu_func      = test_block_cpu_func,
+#ifdef STARPU_USE_CUDA
+	.cuda_func     = test_block_cuda_func,
+#endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_func   = test_block_opencl_func,
+#endif
+	.handle        = &block_handle,
+	.dummy_handle  = &block2_handle,
+	.copy_failed   = 0,
+	.name          = "block_interface"
+};
+
+static int block[NX*NY*NZ];
+static int block2[NX*NY*NZ];
+
+static void
+register_data(void)
+{
+	/* Initializing data */
+	int val = 0;
+	int i, j, k;
+	for (k = 0; k < NZ; k++)
+		for (j = 0; j < NY; j++)
+			for (i = 0; i < NX; i++)
+                                block[(k*NX*NY)+(j*NX)+i] = val++;
+
+	/* Registering data */
+	starpu_block_data_register(&block_handle,
+                                    0,
+                                    (uintptr_t)block,
+				    NX,
+				    NX * NY,
+				    NX,
+				    NY,
+				    NZ,
+				    sizeof(block[0]));
+	starpu_block_data_register(&block2_handle,
+                                    0,
+                                    (uintptr_t)block2,
+				    NX,
+				    NX * NY,
+				    NX,
+				    NY,
+				    NZ,
+				    sizeof(block2[0]));
+}
+
+static void
+unregister_data(void)
+{
+	starpu_data_unregister(block_handle);
+	starpu_data_unregister(block2_handle);
+}
+
+static void test_block_cpu_func(void *buffers[], void *args)
+{
+	int factor = *(int*)args;
+	int nx = STARPU_BLOCK_GET_NX(buffers[0]);
+	int ny = STARPU_BLOCK_GET_NY(buffers[0]);
+	int nz = STARPU_BLOCK_GET_NZ(buffers[0]);
+        unsigned ldy = STARPU_BLOCK_GET_LDY(buffers[0]);
+        unsigned ldz = STARPU_BLOCK_GET_LDZ(buffers[0]);
+	int *block = (int *) STARPU_BLOCK_GET_PTR(buffers[0]);
+	unsigned int i, j, k;
+	int val = 0;
+	block_config.copy_failed = 0;
+	for (k = 0; k < nz; k++)
+	{
+		for (j = 0; j < ny; j++)
+		{
+			for (i = 0; i < nx; i++)
+			{
+                                if (block[(k*ldz)+(j*ldy)+i] != factor * val)
+				{
+					block_config.copy_failed = 1;
+					return;
+				}
+				else
+				{
+					block[(k*ldz)+(j*ldy)+i] *= -1;
+					val++;
+				}
+			}
+		}
+	}
+}
+
+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(&block_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/block/block_opencl.c

@@ -0,0 +1,116 @@
+/* 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/block/block_opencl_kernel.cl"
+extern struct test_config block_config;
+static struct starpu_opencl_program opencl_program;
+
+void
+test_block_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);
+
+	int nx = STARPU_BLOCK_GET_NX(buffers[0]);
+	int ny = STARPU_BLOCK_GET_NY(buffers[0]);
+	int nz = STARPU_BLOCK_GET_NZ(buffers[0]);
+        unsigned ldy = STARPU_BLOCK_GET_LDY(buffers[0]);
+        unsigned ldz = STARPU_BLOCK_GET_LDZ(buffers[0]);
+	cl_mem block = (cl_mem) STARPU_BLOCK_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), &block_config.copy_failed, &err);
+
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+
+	err = starpu_opencl_load_kernel(&kernel,
+					&queue,
+					&opencl_program,
+					"block_opencl",
+					devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	int nargs;
+	nargs = starpu_opencl_set_kernel_args(&err, &kernel,
+					      sizeof(block), &block,
+					      sizeof(nx), &nx,
+					      sizeof(ny), &ny,
+					      sizeof(nz), &nz,
+					      sizeof(ldy), &ldy,
+					      sizeof(ldz), &ldz,
+					      sizeof(factor), &factor,
+					      sizeof(fail), &fail,
+					      0);
+
+	if (nargs != 8)
+	{
+		fprintf(stderr, "Failed to set argument #%d\n", nargs);
+		STARPU_OPENCL_REPORT_ERROR(err);
+	}
+			
+	{
+		size_t global = nx * ny * nz;
+		err = clEnqueueNDRangeKernel(queue,
+					     kernel,
+					     1,
+					     NULL,
+					     &global,
+					     NULL,
+					     0,
+					     NULL,
+					     &event);
+
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	err = clEnqueueReadBuffer(queue,
+				  fail,
+				  CL_TRUE,
+				  0, 
+				  sizeof(int),
+				  &block_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/block/block_opencl_kernel.cl

@@ -0,0 +1,46 @@
+/* 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 block_opencl(__global int *block,
+			   int nx, int ny, int nz,
+			   int ldy, int ldz,
+			   int factor, __global int *err)
+{
+        const int id = get_global_id(0);
+	if (id > 0)
+		return;
+
+	unsigned int i, j, k;
+	int val = 0;
+	for (k = 0; k < nz; k++)
+	{
+		for (j = 0; j < ny; j++)
+		{
+			for (i = 0; i < nx; i++)
+			{
+                                if (block[(k*ldz)+(j*ldy)+i] != factor * val)
+				{
+					*err = 1;
+					return;
+				}
+				else
+				{
+					block[(k*ldz)+(j*ldy)+i] *= -1;
+					val++;
+				}
+			}
+		}
+	}
+}