Add 2d shadow filters

doc/chapters/basic-api.texi

@@ -1160,12 +1160,36 @@ matrices. If nparts does not divide x, the last submatrix contains the
 remainder.
 @end deftypefun
 
+@deftypefun void starpu_block_shadow_filter_func (void *@var{father_interface}, void *@var{child_interface}, {struct starpu_data_filter} *@var{f}, unsigned @var{id}, unsigned @var{nparts})
+This partitions a dense Matrix along the x dimension, with a shadow border
+@code{filter_arg_ptr}, thus getting (y,((x-2*shadow)/nparts+2*shadow)
+matrices. If nparts does not divide x-2*shadow, the last submatrix contains the
+remainder.
+
+IMPORTANT: This can only be used for read-only access, as no coherency is
+enforced for the shadowed parts.
+
+A usage example is available in examples/filters/shadow2d.c
+@end deftypefun
+
 @deftypefun void starpu_vertical_block_filter_func (void *@var{father_interface}, void *@var{child_interface}, {struct starpu_data_filter} *@var{f}, unsigned @var{id}, unsigned @var{nparts})
 This partitions a dense Matrix along the y dimension, thus getting (y/nparts),x)
 matrices. If nparts does not divide y, the last submatrix contains the
 remainder.
 @end deftypefun
 
+@deftypefun void starpu_vertical_block_shadow_filter_func (void *@var{father_interface}, void *@var{child_interface}, {struct starpu_data_filter} *@var{f}, unsigned @var{id}, unsigned @var{nparts})
+This partitions a dense Matrix along the y dimension, with a shadow border
+@code{filter_arg_ptr}, thus getting ((y-2*shadow)/nparts)+2*shadow,x)
+matrices. If nparts does not divide y-2*shadow, the last submatrix contains the
+remainder.
+
+IMPORTANT: This can only be used for read-only access, as no coherency is
+enforced for the shadowed parts.
+
+A usage example is available in examples/filters/shadow2d.c
+@end deftypefun
+
 @node Partitioning Vector Data
 @subsubsection Partitioning Vector Data
 
@@ -1178,7 +1202,7 @@ vector represented by @var{father_interface} once partitioned in
 @deftypefun void starpu_block_shadow_filter_func_vector (void *@var{father_interface}, void *@var{child_interface}, {struct starpu_data_filter} *@var{f}, unsigned @var{id}, unsigned @var{nparts})
 Return in @code{*@var{child_interface}} the @var{id}th element of the
 vector represented by @var{father_interface} once partitioned in
-@var{nparts} chunks of equal size with a shadow border @code{filter_arg_ptr}
+@var{nparts} chunks of equal size with a shadow border @code{filter_arg_ptr}, thus getting a vector of size (n-2*shadow)/nparts+2*shadow 
 
 The @code{filter_arg_ptr} field must be the shadow size casted into @code{void*}.
 

examples/Makefile.am

@@ -180,6 +180,7 @@ examplebin_PROGRAMS +=				\
 	filters/fblock				\
 	filters/fmatrix				\
 	filters/shadow				\
+	filters/shadow2d			\
 	tag_example/tag_example			\
 	tag_example/tag_example2		\
 	tag_example/tag_example3		\

examples/filters/shadow2d.c

@@ -0,0 +1,292 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012  Université de Bordeaux 1
+ * Copyright (C) 2010, 2011, 2012  Centre National de la Recherche Scientifique
+ *
+ * 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.
+ */
+
+/*
+ * This examplifies the use of the matrix shadow filters: a source "matrix" of
+ * NX*NY elements (plus 2*NX*SHADOWX+2*NY*SHADOWY+4*SHADOWX*SHADOWY wrap-around
+ * elements) is partitioned into matrices with some shadowing, and these are
+ * copied into a destination "matrix2" of
+ * NRPARTSX*NPARTSY*((NX/NPARTSX+2*SHADOWX)*(NY/NPARTSY+2*SHADOWY)) elements,
+ * partitioned in the traditionnal way, thus showing how shadowing shows up.
+ *
+ * For instance, with NX=NY=8, SHADOWX=SHADOWY=1, and NPARTSX=NPARTSY=4:
+ *
+ * matrix
+ * 0123456789
+ * 1234567890
+ * 2345678901
+ * 3456789012
+ * 4567890123
+ * 5678901234
+ * 6789012345
+ * 7890123456
+ * 8901234567
+ * 9012345678
+ *
+ * is partitioned into 4*4 pieces:
+ *
+ * 0123 2345 4567 6789
+ * 1234 3456 5678 7890
+ * 2345 4567 6789 8901
+ * 3456 5678 7890 9012
+ *
+ * 2345 4567 6789 8901
+ * 3456 5678 7890 9012
+ * 4567 6789 8901 0123
+ * 5678 7890 9012 1234
+ *
+ * 4567 6789 8901 0123
+ * 5678 7890 9012 1234
+ * 6789 8901 0123 2345
+ * 7890 9012 1234 3456
+ *
+ * 6789 8901 0123 2345
+ * 7890 9012 1234 3456
+ * 8901 0123 2345 4567
+ * 9012 1234 3456 5678
+ *
+ * which are copied into the 4*4 destination subparts of matrix2, thus getting in
+ * the end:
+ *
+ * 0123234545676789
+ * 1234345656787890
+ * 2345456767898901
+ * 3456567878909012
+ * 2345456767898901
+ * 3456567878909012
+ * 4567678989010123
+ * 5678789090121234
+ * 4567678989010123
+ * 5678789090121234
+ * 6789890101232345
+ * 7890901212343456
+ * 6789890101232345
+ * 7890901212343456
+ * 8901012323454567
+ * 9012123434565678
+ */
+
+#include <starpu.h>
+#include <starpu_cuda.h>
+
+/* Shadow width */
+#define SHADOWX 3
+#define SHADOWY 2
+#define NX    20
+#define NY    30
+#define PARTSX 2
+#define PARTSY 3
+
+#define FPRINTF(ofile, fmt, args ...) do { if (!getenv("STARPU_SSILENT")) {fprintf(ofile, fmt, ##args); }} while(0)
+
+void cpu_func(void *buffers[], void *cl_arg)
+{
+        /* length of the shadowed source matrix */
+        unsigned ld = STARPU_MATRIX_GET_LD(buffers[0]);
+        unsigned n = STARPU_MATRIX_GET_NX(buffers[0]);
+        unsigned m = STARPU_MATRIX_GET_NY(buffers[0]);
+        /* local copy of the shadowed source matrix pointer */
+        int *val = (int *)STARPU_MATRIX_GET_PTR(buffers[0]);
+
+        /* length of the destination matrix */
+        unsigned ld2 = STARPU_MATRIX_GET_LD(buffers[1]);
+        unsigned n2 = STARPU_MATRIX_GET_NX(buffers[1]);
+        unsigned m2 = STARPU_MATRIX_GET_NY(buffers[1]);
+        /* local copy of the destination matrix pointer */
+        int *val2 = (int *)STARPU_MATRIX_GET_PTR(buffers[1]);
+
+	unsigned i, j;
+
+	/* If things go right, sizes should match */
+	STARPU_ASSERT(n == n2);
+	STARPU_ASSERT(m == m2);
+	for (j = 0; j < m; j++)
+		for (i = 0; i < n; i++)
+			val2[j*ld2+i] = val[j*ld+i];
+}
+
+#ifdef STARPU_USE_CUDA
+void cuda_func(void *buffers[], void *cl_arg)
+{
+        /* length of the shadowed source matrix */
+        unsigned ld = STARPU_MATRIX_GET_LD(buffers[0]);
+        unsigned n = STARPU_MATRIX_GET_NX(buffers[0]);
+        unsigned m = STARPU_MATRIX_GET_NY(buffers[0]);
+        /* local copy of the shadowed source matrix pointer */
+        int *val = (int *)STARPU_MATRIX_GET_PTR(buffers[0]);
+
+        /* length of the destination matrix */
+        unsigned ld2 = STARPU_MATRIX_GET_LD(buffers[1]);
+        unsigned n2 = STARPU_MATRIX_GET_NX(buffers[1]);
+        unsigned m2 = STARPU_MATRIX_GET_NY(buffers[1]);
+        /* local copy of the destination matrix pointer */
+        int *val2 = (int *)STARPU_MATRIX_GET_PTR(buffers[1]);
+
+	/* If things go right, sizes should match */
+	STARPU_ASSERT(n == n2);
+	STARPU_ASSERT(m == m2);
+	cudaMemcpy2D(val2, ld2*sizeof(*val2), val, ld*sizeof(*val), n*sizeof(*val), m, cudaMemcpyDeviceToDevice);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}
+#endif
+
+int main(int argc, char **argv)
+{
+	unsigned i, j, k, l;
+        int matrix[NY + 2*SHADOWY][NX + 2*SHADOWX];
+        int matrix2[NY + PARTSY*2*SHADOWY][NX + PARTSX*2*SHADOWX];
+	starpu_data_handle_t handle, handle2;
+	int ret;
+
+        struct starpu_codelet cl =
+	{
+                .where = STARPU_CPU
+#ifdef STARPU_USE_CUDA
+			|STARPU_CUDA
+#endif
+			,
+                .cpu_funcs = {cpu_func, NULL},
+#ifdef STARPU_USE_CUDA
+                .cuda_funcs = {cuda_func, NULL},
+#endif
+                .nbuffers = 2,
+		.modes = {STARPU_R, STARPU_W}
+        };
+
+	memset(matrix, -1, sizeof(matrix));
+        for(i=1 ; i<=NX ; i++)
+		for(j=1 ; j<=NY ; j++)
+			matrix[SHADOWY+j-1][SHADOWX+i-1] = i+j;
+
+	/* Copy borders */
+	for(i=0 ; i<SHADOWX ; i++)
+		for (j = SHADOWY ; j<SHADOWY+NY ; j++) {
+			matrix[j][i] = matrix[j][i+NX];
+			matrix[j][SHADOWX+NX+i] = matrix[j][SHADOWX+i];
+		}
+	for(i=SHADOWX ; i<SHADOWX+NX ; i++)
+		for(j=0 ; j<SHADOWY ; j++) {
+			matrix[j][i] = matrix[j+NY][i];
+			matrix[SHADOWY+NY+j][i] = matrix[SHADOWY+j][i];
+		}
+	/* Copy corners */
+	for(i=0 ; i<SHADOWX ; i++)
+		for(j=0 ; j<SHADOWY ; j++) {
+			matrix[j][i] = matrix[j+NY][i+NX];
+			matrix[j][SHADOWX+NX+i] = matrix[j+NY][SHADOWX+i];
+			matrix[SHADOWY+NY+j][i] = matrix[SHADOWY+j][i+NX];
+			matrix[SHADOWY+NY+j][SHADOWX+NX+i] = matrix[SHADOWY+j][SHADOWX+i];
+		}
+
+        FPRINTF(stderr,"IN  Matrix:\n");
+	for(j=0 ; j<NY + 2*SHADOWY ; j++)
+	{
+		for(i=0 ; i<NX + 2*SHADOWX ; i++)
+			FPRINTF(stderr, "%5d ", matrix[j][i]);
+		FPRINTF(stderr,"\n");
+	}
+        FPRINTF(stderr,"\n");
+
+	ret = starpu_init(NULL);
+	if (ret == -ENODEV)
+		exit(77);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_init");
+
+	/* Declare source matrix to StarPU */
+	starpu_matrix_data_register(&handle, 0, (uintptr_t)matrix, NX + 2*SHADOWX, NX + 2*SHADOWX, NY + 2*SHADOWY, sizeof(matrix[0][0]));
+
+	/* Declare destination matrix to StarPU */
+	starpu_matrix_data_register(&handle2, 0, (uintptr_t)matrix2, NX + PARTSX*2*SHADOWX, NX + PARTSX*2*SHADOWX, NY + PARTSY*2*SHADOWY, sizeof(matrix2[0][0]));
+
+        /* Partition the source matrix in PARTSY*PARTSX sub-matrices with shadows */
+	/* NOTE: the resulting handles should onlx be used in read-onlx mode,
+	 * as StarPU will not know how the overlapping parts would have to be
+	 * combined. */
+	struct starpu_data_filter fy =
+	{
+		.filter_func = starpu_vertical_block_shadow_filter_func,
+		.nchildren = PARTSY,
+		.filter_arg_ptr = (void*)(uintptr_t) SHADOWY /* Shadow width */
+	};
+	struct starpu_data_filter fx =
+	{
+		.filter_func = starpu_block_shadow_filter_func,
+		.nchildren = PARTSX,
+		.filter_arg_ptr = (void*)(uintptr_t) SHADOWX /* Shadow width */
+	};
+	starpu_data_map_filters(handle, 2, &fy, &fx);
+
+        /* Partition the destination matrix in PARTSY*PARTSX sub-matrices */
+	struct starpu_data_filter fy2 =
+	{
+		.filter_func = starpu_vertical_block_filter_func,
+		.nchildren = PARTSY,
+	};
+	struct starpu_data_filter fx2 =
+	{
+		.filter_func = starpu_block_filter_func,
+		.nchildren = PARTSX,
+	};
+	starpu_data_map_filters(handle2, 2, &fy2, &fx2);
+
+        /* Submit a task on each sub-vector */
+	for (j=0; j<PARTSY; j++)
+	{
+		for (i=0; i<PARTSX; i++)
+		{
+			starpu_data_handle_t sub_handle = starpu_data_get_sub_data(handle, 2, j, i);
+			starpu_data_handle_t sub_handle2 = starpu_data_get_sub_data(handle2, 2, j, i);
+			struct starpu_task *task = starpu_task_create();
+
+			task->handles[0] = sub_handle;
+			task->handles[1] = sub_handle2;
+			task->cl = &cl;
+			task->synchronous = 1;
+
+			ret = starpu_task_submit(task);
+			if (ret == -ENODEV) goto enodev;
+			STARPU_CHECK_RETURN_VALUE(ret, "starpu_task_submit");
+		}
+	}
+
+	starpu_data_unpartition(handle, 0);
+	starpu_data_unpartition(handle2, 0);
+        starpu_data_unregister(handle);
+        starpu_data_unregister(handle2);
+	starpu_shutdown();
+
+        FPRINTF(stderr,"OUT Matrix:\n");
+	for(j=0 ; j<NY + PARTSY*2*SHADOWY ; j++)
+	{
+		for(i=0 ; i<NX + PARTSX*2*SHADOWX ; i++)
+			FPRINTF(stderr, "%5d ", matrix2[j][i]);
+		FPRINTF(stderr,"\n");
+	}
+        FPRINTF(stderr,"\n");
+	for(i=0 ; i<PARTSX ; i++)
+		for(j=0 ; j<PARTSY ; j++)
+			for (k=0 ; k<NX/PARTSX ; k++)
+				for (l=0 ; l<NY/PARTSY ; l++)
+					STARPU_ASSERT(matrix2[j*(NX/PARTSY+2*SHADOWY)+l][i*(NX/PARTSX+2*SHADOWX)+k] = matrix[i*(NX/PARTSX)+k][j*(NY/PARTSY)+l]);
+
+	return 0;
+
+enodev:
+	FPRINTF(stderr, "WARNING: No one can execute this task\n");
+	starpu_shutdown();
+	return 77;
+}

src/datawizard/interfaces/matrix_filters.c

@@ -1,6 +1,6 @@
 /* StarPU --- Runtime system for heterogeneous multicore architectures.
  *
- * Copyright (C) 2010-2011  Université de Bordeaux 1
+ * Copyright (C) 2010-2012  Université de Bordeaux 1
  * Copyright (C) 2010  Mehdi Juhoor <mjuhoor@gmail.com>
  * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
  *
@@ -56,6 +56,45 @@ void starpu_block_filter_func(void *father_interface, void *child_interface, STA
 	}
 }
 
+/*
+ * an example of a dummy partition function : blocks ...
+ */
+void starpu_block_shadow_filter_func(void *father_interface, void *child_interface, STARPU_ATTRIBUTE_UNUSED struct starpu_data_filter *f, unsigned id, unsigned nchunks)
+{
+	struct starpu_matrix_interface *matrix_father = (struct starpu_matrix_interface *) father_interface;
+	struct starpu_matrix_interface *matrix_child = (struct starpu_matrix_interface *) child_interface;
+
+	uintptr_t shadow_size = (uintptr_t) f->filter_arg_ptr;
+
+	/* actual number of elements */
+	uint32_t nx = matrix_father->nx - 2 * shadow_size;
+	uint32_t ny = matrix_father->ny;
+	size_t elemsize = matrix_father->elemsize;
+
+	STARPU_ASSERT(nchunks <= nx);
+
+	size_t chunk_size = ((size_t)nx + nchunks - 1)/nchunks;
+	size_t offset = (size_t)id*chunk_size*elemsize;
+
+	uint32_t child_nx =
+	  STARPU_MIN(chunk_size, (size_t)nx - (size_t)id*chunk_size) + 2 * shadow_size;
+
+	/* update the child's interface */
+	matrix_child->nx = child_nx;
+	matrix_child->ny = ny;
+	matrix_child->elemsize = elemsize;
+
+	/* is the information on this node valid ? */
+	if (matrix_father->dev_handle)
+	{
+		if (matrix_father->ptr)
+			matrix_child->ptr = matrix_father->ptr + offset;
+		matrix_child->ld = matrix_father->ld;
+		matrix_child->dev_handle = matrix_father->dev_handle;
+		matrix_child->offset = matrix_father->offset + offset;
+	}
+}
+
 void starpu_vertical_block_filter_func(void *father_interface, void *child_interface, STARPU_ATTRIBUTE_UNUSED struct starpu_data_filter *f, unsigned id, unsigned nchunks)
 {
         struct starpu_matrix_interface *matrix_father = (struct starpu_matrix_interface *) father_interface;
@@ -86,3 +125,37 @@ void starpu_vertical_block_filter_func(void *father_interface, void *child_inter
 		matrix_child->offset = matrix_father->offset + offset;
 	}
 }
+
+void starpu_vertical_block_shadow_filter_func(void *father_interface, void *child_interface, STARPU_ATTRIBUTE_UNUSED struct starpu_data_filter *f, unsigned id, unsigned nchunks)
+{
+        struct starpu_matrix_interface *matrix_father = (struct starpu_matrix_interface *) father_interface;
+        struct starpu_matrix_interface *matrix_child = (struct starpu_matrix_interface *) child_interface;
+
+	uintptr_t shadow_size = (uintptr_t) f->filter_arg_ptr;
+
+	uint32_t nx = matrix_father->nx;
+	/* actual number of elements */
+	uint32_t ny = matrix_father->ny - 2 * shadow_size;
+	size_t elemsize = matrix_father->elemsize;
+
+	STARPU_ASSERT(nchunks <= ny);
+
+	size_t chunk_size = ((size_t)ny + nchunks - 1)/nchunks;
+	size_t child_ny =
+	  STARPU_MIN(chunk_size, (size_t)ny - (size_t)id*chunk_size) + 2 * shadow_size;
+
+	matrix_child->nx = nx;
+	matrix_child->ny = child_ny;
+	matrix_child->elemsize = elemsize;
+
+	/* is the information on this node valid ? */
+	if (matrix_father->dev_handle)
+	{
+		size_t offset = (size_t)id*chunk_size*matrix_father->ld*elemsize;
+		if (matrix_father->ptr)
+			matrix_child->ptr = matrix_father->ptr + offset;
+		matrix_child->ld = matrix_father->ld;
+		matrix_child->dev_handle = matrix_father->dev_handle;
+		matrix_child->offset = matrix_father->offset + offset;
+	}
+}