Make tests/perfmodels/regression_based and tests/perfmodels/non_linear_regression_based work with OpenCL.

tests/Makefile.am

@@ -40,7 +40,8 @@ EXTRA_DIST =					\
 	datawizard/interfaces/multiformat/multiformat_conversion_codelets_kernel.cl \
 	datawizard/interfaces/multiformat/advanced/generic.h \
 	datawizard/interfaces/csr/csr_opencl_kernel.cl \
-	datawizard/interfaces/block/block_opencl_kernel.cl
+	datawizard/interfaces/block/block_opencl_kernel.cl \
+	perfmodels/opencl_memset_kernel.cl
 
 CLEANFILES = 					\
 	*.gcno *.gcda *.linkinfo		\
@@ -521,5 +522,26 @@ datawizard_interfaces_void_void_interface_SOURCES=\
 	datawizard/interfaces/test_interfaces.c        \
 	datawizard/interfaces/void/void_interface.c
 
+
+perfmodels_regression_based_SOURCES=\
+	perfmodels/regression_based.c
+
+if STARPU_USE_OPENCL
+perfmodels_regression_based_SOURCES+=\
+	perfmodels/opencl_memset.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	perfmodels/opencl_memset_kernel.cl
+endif
+
+perfmodels_non_linear_regression_based_SOURCES=\
+	perfmodels/non_linear_regression_based.c
+
+if STARPU_USE_OPENCL
+perfmodels_non_linear_regression_based_SOURCES+=\
+	perfmodels/opencl_memset.c
+nobase_STARPU_OPENCL_DATA_DATA += \
+	perfmodels/opencl_memset_kernel.cl
+endif
+
 showcheck:
 	-cat $(TEST_LOGS) /dev/null

tests/perfmodels/non_linear_regression_based.c

@@ -2,6 +2,7 @@
  *
  * Copyright (C) 2011  Université de Bordeaux 1
  * Copyright (C) 2012  Centre National de la Recherche Scientifique
+ * Copyright (C) 2012 inria
  *
  * 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
@@ -17,6 +18,9 @@
 
 #include <config.h>
 #include <starpu.h>
+#ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
+#endif
 #include "../helper.h"
 
 #ifdef STARPU_USE_CUDA
@@ -48,12 +52,18 @@ static struct starpu_perfmodel model =
 	.symbol = "non_linear_memset_regression_based"
 };
 
+#ifdef STARPU_USE_OPENCL
+extern void memset_opencl(void *buffers[], void *args);
+#endif
+
 static struct starpu_codelet memset_cl =
 {
-	.where = STARPU_CUDA|STARPU_CPU,
 #ifdef STARPU_USE_CUDA
 	.cuda_funcs = {memset_cuda, NULL},
 #endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_funcs = {memset_opencl, NULL},
+#endif
 	.cpu_funcs = {memset_cpu, NULL},
 	.model = &model,
 	.nbuffers = 1,
@@ -84,6 +94,10 @@ static void test_memset(int nelems)
 	starpu_data_unregister(handle);
 }
 
+#ifdef STARPU_USE_OPENCL
+struct starpu_opencl_program opencl_program;
+#endif
+
 int main(int argc, char **argv)
 {
 	int ret;
@@ -98,6 +112,12 @@ int main(int argc, char **argv)
 	if (ret == -ENODEV) return STARPU_TEST_SKIPPED;
 	STARPU_CHECK_RETURN_VALUE(ret, "starpu_init");
 
+#ifdef STARPU_USE_OPENCL
+	ret = starpu_opencl_load_opencl_from_file("tests/perfmodels/opencl_memset_kernel.cl",
+						  &opencl_program, NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+#endif
+
 	int slog;
 	for (slog = 8; slog < 25; slog++)
 	{
@@ -105,6 +125,10 @@ int main(int argc, char **argv)
 		test_memset(size);
 	}
 
+#ifdef STARPU_USE_OPENCL
+        starpu_opencl_unload_opencl(&opencl_program);
+#endif
+
 	starpu_shutdown();
 
 	return EXIT_SUCCESS;

tests/perfmodels/opencl_memset.c

@@ -0,0 +1,69 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 inria
+ *
+ * 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>
+
+extern struct starpu_opencl_program opencl_program;
+
+void memset_opencl(void *buffers[], void *args)
+{
+	(void) *args;
+	int id, devid;
+        cl_int err;
+	cl_kernel kernel;
+	cl_command_queue queue;
+	cl_event event;
+
+	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	cl_mem val = (cl_mem)STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
+
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_program, "_memset_opencl", devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 0, sizeof(val), &val);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		size_t global=n;
+		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);
+	}
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+}

tests/perfmodels/opencl_memset_kernel.cl

@@ -0,0 +1,22 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 inria
+ *
+ * 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 _memset_opencl(__global int *val, int nx)
+{
+        const int i = get_global_id(0);
+        if (i < nx)
+                val[i] = 42;
+}

tests/perfmodels/regression_based.c

@@ -2,6 +2,7 @@
  *
  * Copyright (C) 2011-2012  Université de Bordeaux 1
  * Copyright (C) 2011  Télécom-SudParis
+ * Copyright (C) 2012 inria
  *
  * 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
@@ -17,6 +18,9 @@
 
 #include <config.h>
 #include <starpu.h>
+#ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
+#endif
 #include "../helper.h"
 
 #ifdef STARPU_USE_CUDA
@@ -32,6 +36,10 @@ static void memset_cuda(void *descr[], void *arg)
 }
 #endif
 
+#ifdef STARPU_USE_OPENCL
+extern void memset_opencl(void *buffers[], void *args);
+#endif
+
 static void memset_cpu(void *descr[], void *arg)
 {
 	STARPU_SKIP_IF_VALGRIND;
@@ -56,10 +64,12 @@ static struct starpu_perfmodel nl_model =
 
 static struct starpu_codelet memset_cl =
 {
-	.where = STARPU_CUDA|STARPU_CPU,
 #ifdef STARPU_USE_CUDA
 	.cuda_funcs = {memset_cuda, NULL},
 #endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_funcs = {memset_opencl, NULL},
+#endif
 	.cpu_funcs = {memset_cpu, NULL},
 	.model = &model,
 	.nbuffers = 1,
@@ -68,10 +78,12 @@ static struct starpu_codelet memset_cl =
 
 static struct starpu_codelet nl_memset_cl =
 {
-	.where = STARPU_CUDA|STARPU_CPU,
 #ifdef STARPU_USE_CUDA
 	.cuda_funcs = {memset_cuda, NULL},
 #endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_funcs = {memset_opencl, NULL},
+#endif
 	.cpu_funcs = {memset_cpu, NULL},
 	.model = &nl_model,
 	.nbuffers = 1,
@@ -118,6 +130,10 @@ static void show_task_perfs(int size, struct starpu_task *task)
 	}
 }
 
+#ifdef STARPU_USE_OPENCL
+struct starpu_opencl_program opencl_program;
+#endif
+
 int main(int argc, char **argv)
 {
 	struct starpu_conf conf;
@@ -133,6 +149,12 @@ int main(int argc, char **argv)
 	if (ret == -ENODEV) return STARPU_TEST_SKIPPED;
 	STARPU_CHECK_RETURN_VALUE(ret, "starpu_init");
 
+#ifdef STARPU_USE_OPENCL
+	ret = starpu_opencl_load_opencl_from_file("tests/perfmodels/opencl_memset_kernel.cl",
+						  &opencl_program, NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+#endif
+
 	int size;
 	for (size = 1024; size < 16777216; size *= 2)
 	{
@@ -167,6 +189,9 @@ int main(int argc, char **argv)
 
 	starpu_data_unregister(handle);
 
+#ifdef STARPU_USE_OPENCL
+        starpu_opencl_unload_opencl(&opencl_program);
+#endif
 	starpu_shutdown();
 
 	return EXIT_SUCCESS;