Filter block example: add OpenCL implementation

examples/Makefile.am

@@ -179,6 +179,13 @@ if STARPU_USE_CUDA
 filters_fblock_SOURCES +=			\
 	filters/fblock_cuda.cu
 endif
+if STARPU_USE_OPENCL
+filters_fblock_SOURCES +=			\
+	filters/fblock_opencl.c
+endif
+
+nobase_STARPU_OPENCL_DATA_DATA += \
+	filters/fblock_opencl_codelet.cl
 
 filters_fmatrix_SOURCES =			\
 	filters/fmatrix.c

examples/filters/fblock.c

@@ -1,6 +1,6 @@
 /*
  * StarPU
- * Copyright (C) INRIA 2008-2009 (see AUTHORS file)
+ * Copyright (C) INRIA 2008-2010 (see AUTHORS file)
  *
  * This program is free software; you can redistribute it and/or modify
  * it under the terms of the GNU Lesser General Public License as published by
@@ -22,18 +22,13 @@
 #define NZ    3
 #define PARTS 2
 
-void opencl_func(void *buffers[], void *cl_arg)
-{
-        return;
-}
-
 extern void cuda_func(void *buffers[], void *cl_arg);
+extern void opencl_func(void *buffers[], void *cl_arg);
 
 void cpu_func(void *buffers[], void *cl_arg)
 {
         unsigned i, j, k;
         int *factor = cl_arg;
-
 	int *block = (int *)STARPU_GET_BLOCK_PTR(buffers[0]);
 	int nx = (int)STARPU_GET_BLOCK_NX(buffers[0]);
 	int ny = (int)STARPU_GET_BLOCK_NY(buffers[0]);
@@ -49,14 +44,14 @@ void cpu_func(void *buffers[], void *cl_arg)
         }
 }
 
-void print_block(int *block, int nx, int ny, int nz)
+void print_block(int *block, int nx, int ny, int nz, unsigned ldy, unsigned ldz)
 {
         int i, j, k;
-
+        fprintf(stderr, "block=%p nx=%d ny=%d nz=%d ldy=%d ldz=%d\n", block, nx, ny, nz, ldy, ldz);
         for(k=0 ; k<nz ; k++) {
                 for(j=0 ; j<ny ; j++) {
                         for(i=0 ; i<nx ; i++) {
-                                fprintf(stderr, "%2d ", block[(k*nx*ny)+(j*nx)+i]);
+                                fprintf(stderr, "%2d ", block[(k*ldz)+(j*ldy)+i]);
                         }
                         fprintf(stderr,"\n");
                 }
@@ -65,6 +60,22 @@ void print_block(int *block, int nx, int ny, int nz)
         fprintf(stderr,"\n");
 }
 
+void print_data(starpu_data_handle block_handle)
+{
+	int *block = starpu_block_get_local_ptr(block_handle);
+	int nx = starpu_block_get_nx(block_handle);
+	int ny = starpu_block_get_ny(block_handle);
+	int nz = starpu_block_get_nz(block_handle);
+	unsigned ldy = starpu_block_get_local_ldy(block_handle);
+	unsigned ldz = starpu_block_get_local_ldz(block_handle);
+
+        print_block(block, nx, ny, nz, ldy, ldz);
+}
+
+#ifdef STARPU_USE_OPENCL
+struct starpu_opencl_program codelet;
+#endif
+
 int main(int argc, char **argv)
 {
         int *block,n=0;
@@ -79,21 +90,26 @@ int main(int argc, char **argv)
                         }
                 }
         }
-        fprintf(stderr, "IN  Block\n");
-        print_block(block, NX, NY, NZ);
 
 	starpu_data_handle handle;
 	starpu_codelet cl =
 	{
-		.where = STARPU_CPU|STARPU_CUDA,
+                .where = STARPU_CPU|STARPU_CUDA|STARPU_OPENCL,
                 .cpu_func = cpu_func,
                 .cuda_func = cuda_func,
+                .opencl_func = opencl_func,
 		.nbuffers = 1
 	};
         starpu_init(NULL);
 
+#ifdef STARPU_USE_OPENCL
+        starpu_opencl_load_opencl_from_file("examples/filters/fblock_opencl_codelet.cl", &codelet);
+#endif
+
         /* Declare data to StarPU */
         starpu_block_data_register(&handle, 0, (uintptr_t)block, NX, NX*NY, NX, NY, NZ, sizeof(int));
+        fprintf(stderr, "IN  Block\n");
+        print_data(handle);
 
         /* Partition the block in PARTS sub-blocks */
 	struct starpu_data_filter f =
@@ -107,6 +123,13 @@ int main(int argc, char **argv)
 
         fprintf(stderr,"Nb of partitions : %d\n",starpu_data_get_nb_children(handle));
 
+        for(i=0 ; i<starpu_data_get_nb_children(handle) ; i++)
+        {
+                starpu_data_handle sblock = starpu_data_get_sub_data(handle, 1, i);
+                fprintf(stderr, "Sub block %d\n", i);
+                print_data(sblock);
+        }
+
         /* Submit a task on each sub-block */
         for(i=0 ; i<starpu_data_get_nb_children(handle) ; i++)
         {
@@ -130,10 +153,12 @@ int main(int argc, char **argv)
 
         /* Unpartition the data, unregister it from StarPU and shutdown */
         starpu_data_unpartition(handle, 0);
+        print_data(handle);
         starpu_data_unregister(handle);
-	starpu_shutdown();
 
         /* Print result block */
         fprintf(stderr, "OUT Block\n");
-        print_block(block, NX, NY, NZ);
+        print_block(block, NX, NY, NZ, NX, NX*NY);
+
+	starpu_shutdown();
 }

examples/filters/fblock_opencl.c

@@ -0,0 +1,62 @@
+/*
+ * StarPU
+ * Copyright (C) INRIA 2008-2010 (see AUTHORS file)
+ *
+ * This program 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.
+ *
+ * This program 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>
+
+extern struct starpu_opencl_program codelet;
+
+void opencl_func(void *buffers[], void *cl_arg)
+{
+	int id, devid, err;
+	cl_kernel kernel;
+	cl_command_queue queue;
+
+        int *factor = cl_arg;
+	int *block = (int *)STARPU_GET_BLOCK_PTR(buffers[0]);
+	int nx = (int)STARPU_GET_BLOCK_NX(buffers[0]);
+	int ny = (int)STARPU_GET_BLOCK_NY(buffers[0]);
+	int nz = (int)STARPU_GET_BLOCK_NZ(buffers[0]);
+        unsigned ldy = STARPU_GET_BLOCK_LDY(buffers[0]);
+        unsigned ldz = STARPU_GET_BLOCK_LDZ(buffers[0]);
+
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+
+	err = starpu_opencl_load_kernel(&kernel, &queue, &codelet, "fblock_opencl", devid);
+	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = 0;
+	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &block);
+	err = clSetKernelArg(kernel, 1, sizeof(nx), &nx);
+	err = clSetKernelArg(kernel, 2, sizeof(ny), &ny);
+	err = clSetKernelArg(kernel, 3, sizeof(nz), &nz);
+	err = clSetKernelArg(kernel, 4, sizeof(ldy), &ldy);
+	err = clSetKernelArg(kernel, 5, sizeof(ldz), &ldz);
+	err |= clSetKernelArg(kernel, 6, sizeof(*factor), factor);
+	if (err) STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		size_t global=nx*ny*nz;
+		err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
+		if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	clFinish(queue);
+
+	starpu_opencl_release_kernel(kernel);
+}
+

examples/filters/fblock_opencl_codelet.cl

@@ -0,0 +1,30 @@
+/*
+ * StarPU
+ * Copyright (C) INRIA 2008-2010 (see AUTHORS file)
+ *
+ * This program 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.
+ *
+ * This program 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 fblock_opencl(__global int* block, int nx, int ny, int nz, unsigned ldy, unsigned ldz, int factor)
+{
+        int i, j, k;
+        for(k=0; k<nz ; k++) {
+                for(j=0; j<ny ; j++) {
+                        for(i=0; i<nx ; i++)
+                                block[(k*ldz)+(j*ldy)+i] = factor;
+                }
+        }
+//        const int i = get_global_id(0);
+//        if (i < nx*ny*nz) {
+//                block[i] = factor;//10*(factor+1);
+//        }
+}