Example: update towards new OpenCL API

examples/basic_examples/vector_scal.c

@@ -26,6 +26,7 @@
 #include <stdio.h>
 #include <stdint.h>
 #include <starpu.h>
+#include <starpu_opencl.h>
 
 #define	NX	2048
 
@@ -86,6 +87,10 @@ static starpu_codelet cl = {
 	.nbuffers = 1
 };
 
+#ifdef STARPU_USE_OPENCL
+struct starpu_opencl_codelet codelet;
+#endif
+
 int main(int argc, char **argv)
 {
 	/* We consider a vector of float that is initialized just as any of C
@@ -101,7 +106,7 @@ int main(int argc, char **argv)
 	starpu_init(NULL);
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("examples/basic_examples/vector_scal_opencl_codelet.cl");
+        starpu_opencl_load_opencl_from_file("examples/basic_examples/vector_scal_opencl_codelet.cl", &codelet);
 #endif
 
 	/* Tell StaPU to associate the "vector" vector with the "vector_handle"
@@ -146,6 +151,10 @@ int main(int argc, char **argv)
  	 * monitoring it */
 	starpu_data_unregister(vector_handle);
 
+#ifdef STARPU_USE_OPENCL
+        starpu_opencl_unload_opencl(&codelet);
+#endif
+
 	/* terminate StarPU, no task can be submitted after */
 	starpu_shutdown();
 

examples/basic_examples/vector_scal_opencl.c

@@ -21,6 +21,8 @@
 #include <starpu.h>
 #include <starpu_opencl.h>
 
+extern struct starpu_opencl_codelet codelet;
+
 void scal_opencl_func(void *buffers[], void *_args)
 {
 	float *factor = (float *)_args;
@@ -37,9 +39,7 @@ void scal_opencl_func(void *buffers[], void *_args)
 	id = starpu_worker_get_id();
 	devid = starpu_worker_get_devid(id);
 
-	err = starpu_opencl_load_kernel(&kernel, &queue,
-					"examples/basic_examples/vector_scal_opencl_codelet.cl",
-					"vector_mult_opencl", devid);
+	err = starpu_opencl_load_kernel(&kernel, &queue, &codelet, "vector_mult_opencl", devid);
 	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -49,13 +49,13 @@ void scal_opencl_func(void *buffers[], void *_args)
 	if (err) STARPU_OPENCL_REPORT_ERROR(err);
 
 	{
-		size_t global=1;
-		size_t local=1;
+		size_t global=n;
+		size_t local=n;
 		err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
 		if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 	}
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }

examples/block/block.c

@@ -32,6 +32,7 @@ void cpu_codelet(void *descr[], void *_args)
 }
 
 #ifdef STARPU_USE_OPENCL
+struct starpu_opencl_codelet opencl_code;
 void opencl_codelet(void *descr[], void *_args)
 {
 	cl_kernel kernel;
@@ -46,8 +47,7 @@ void opencl_codelet(void *descr[], void *_args)
         id = starpu_worker_get_id();
         devid = starpu_worker_get_devid(id);
 
-        err = starpu_opencl_load_kernel(&kernel, &queue,
-                                        "examples/block/block_kernel.cl", "block", devid);
+        err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "block", devid);
         if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -67,7 +67,7 @@ void opencl_codelet(void *descr[], void *_args)
 
 	clFinish(queue);
 
-        starpu_opencl_release(kernel);
+        starpu_opencl_release_kernel(kernel);
 }
 #endif
 
@@ -139,7 +139,7 @@ int main(int argc, char **argv)
         ret = execute_on(STARPU_CPU, cpu_codelet, block, nx, ny, nz, 1.0);
         if (!ret) multiplier *= 1.0;
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("examples/block/block_kernel.cl");
+        starpu_opencl_load_opencl_from_file("examples/block/block_kernel.cl", &opencl_code);
         ret = execute_on(STARPU_OPENCL, opencl_codelet, block, nx, ny, nz, 2.0);
         if (!ret) multiplier *= 2.0;
 #endif

examples/incrementer/incrementer.c

@@ -25,7 +25,9 @@ extern void cuda_codelet(void *descr[], __attribute__ ((unused)) void *_args);
 #endif
 
 #ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
 extern void opencl_codelet(void *descr[], __attribute__ ((unused)) void *_args);
+struct starpu_opencl_codelet opencl_code;
 #endif
 
 extern void cuda_codelet_host(float *tab);
@@ -51,7 +53,7 @@ int main(int argc, char **argv)
 			(uintptr_t)&float_array, 4, sizeof(float));
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("examples/incrementer/incrementer_kernels_opencl_codelet.cl");
+        starpu_opencl_load_opencl_from_file("examples/incrementer/incrementer_kernels_opencl_codelet.cl", &opencl_code);
 #endif
 
 	starpu_codelet cl =

examples/incrementer/incrementer_kernels_opencl.c

@@ -18,6 +18,7 @@
 #include <starpu_opencl.h>
 #include <CL/cl.h>
 
+extern struct starpu_opencl_codelet opencl_code;
 void opencl_codelet(void *descr[], void *_args)
 {
 	float *val = (float *)STARPU_GET_VECTOR_PTR(descr[0]);
@@ -28,8 +29,7 @@ void opencl_codelet(void *descr[], void *_args)
         id = starpu_worker_get_id();
         devid = starpu_worker_get_devid(id);
 
-	err = starpu_opencl_load_kernel(&kernel, &queue,
-                                        "examples/incrementer/incrementer_kernels_opencl_codelet.cl", "incrementer", devid);
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "incrementer", devid);
 	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -45,5 +45,5 @@ void opencl_codelet(void *descr[], void *_args)
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }

examples/matvecmult/matvecmult.c

@@ -25,6 +25,7 @@ static int width=20;
 static int height=4;
 
 #ifdef STARPU_USE_OPENCL
+struct starpu_opencl_codelet opencl_code;
 void opencl_codelet(void *descr[], __attribute__ ((unused)) void *_args)
 {
 	cl_kernel kernel;
@@ -37,8 +38,7 @@ void opencl_codelet(void *descr[], __attribute__ ((unused)) void *_args)
         id = starpu_worker_get_id();
         devid = starpu_worker_get_devid(id);
 
-        err = starpu_opencl_load_kernel(&kernel, &queue,
-                                        "examples/matvecmult/matvecmult_kernel.cl", "matVecMult", devid);
+        err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "matVecMult", devid);
         if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -58,7 +58,7 @@ void opencl_codelet(void *descr[], __attribute__ ((unused)) void *_args)
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }
 #endif
 
@@ -152,7 +152,7 @@ int main(int argc, char **argv)
 	starpu_vector_data_register(&mult_handle, 0, (uintptr_t)mult, height, sizeof(float));
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("examples/matvecmult/matvecmult_kernel.cl");
+        starpu_opencl_load_opencl_from_file("examples/matvecmult/matvecmult_kernel.cl", &opencl_code);
 #endif
 
 	cl.where = STARPU_OPENCL;

examples/spmv/dw_spmv.c

@@ -29,6 +29,7 @@ struct timeval end;
 
 #ifdef STARPU_USE_OPENCL
 #include "starpu_opencl.h"
+struct starpu_opencl_codelet opencl_codelet;
 void spmv_kernel_opencl(void *descr[], void *args)
 {
 	cl_kernel kernel;
@@ -51,8 +52,7 @@ void spmv_kernel_opencl(void *descr[], void *args)
         id = starpu_worker_get_id();
         devid = starpu_worker_get_devid(id);
 
-        err = starpu_opencl_load_kernel(&kernel, &queue,
-                                        "examples/spmv/spmv_opencl.cl", "spvm", devid);
+        err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_codelet, "spvm", devid);
         if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -77,7 +77,7 @@ void spmv_kernel_opencl(void *descr[], void *args)
 
 	clFinish(queue);
 
-        starpu_opencl_release(kernel);
+        starpu_opencl_release_kernel(kernel);
 }
 #endif
 
@@ -245,7 +245,7 @@ void call_spmv_codelet_filters(void)
 
 #ifdef STARPU_USE_OPENCL
         {
-                int ret = _starpu_opencl_compile_source_to_opencl("examples/spmv/spmv_opencl.cl");
+                int ret = starpu_opencl_load_opencl_from_file("examples/spmv/spmv_opencl.cl", &opencl_codelet);
                 if (ret)
 		{
 			fprintf(stderr, "Failed to compile OpenCL codelet\n");

examples/variable/variable.c

@@ -24,7 +24,9 @@ extern void cuda_codelet(void *descr[], __attribute__ ((unused)) void *_args);
 #endif
 
 #ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
 extern void opencl_codelet(void *descr[], __attribute__ ((unused)) void *_args);
+struct starpu_opencl_codelet opencl_code;
 #endif
 
 extern void cuda_codelet_host(float *tab);
@@ -51,7 +53,7 @@ int main(int argc, char **argv)
                                       (uintptr_t)&foo, sizeof(float));
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("examples/variable/variable_kernels_opencl_codelet.cl");
+        starpu_opencl_load_opencl_from_file("examples/variable/variable_kernels_opencl_codelet.cl", &opencl_code);
 #endif
 
 	cl.where = STARPU_CPU|STARPU_CUDA|STARPU_OPENCL;

examples/variable/variable_kernels_opencl.c

@@ -17,6 +17,7 @@
 #include <starpu.h>
 #include <starpu_opencl.h>
 
+extern struct starpu_opencl_codelet opencl_code;
 void opencl_codelet(void *descr[], void *_args)
 {
 	float *val = (float *)STARPU_GET_VECTOR_PTR(descr[0]);
@@ -27,8 +28,7 @@ void opencl_codelet(void *descr[], void *_args)
         id = starpu_worker_get_id();
         devid = starpu_worker_get_devid(id);
 
-	err = starpu_opencl_load_kernel(&kernel, &queue,
-                                        "examples/variable/variable_kernels_opencl_codelet.cl", "variable", devid);
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "variable", devid);
 	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -44,5 +44,5 @@ void opencl_codelet(void *descr[], void *_args)
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }

tests/datawizard/sync_and_notify_data.c

@@ -47,8 +47,10 @@ void cuda_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args);
 #endif
 
 #ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
 void opencl_codelet_incA(void *descr[], __attribute__ ((unused)) void *_args);
 void opencl_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args);
+struct starpu_opencl_codelet opencl_code;
 #endif
 
 #define VECTORSIZE	16
@@ -86,7 +88,7 @@ int main(int argc, char **argv)
 #endif
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("tests/datawizard/sync_and_notify_data_opencl_codelet.cl");
+        starpu_opencl_load_opencl_from_file("tests/datawizard/sync_and_notify_data_opencl_codelet.cl", &opencl_code);
 #endif
 
         starpu_vector_data_register(&v_handle, 0, (uintptr_t)v, VECTORSIZE, sizeof(unsigned));

tests/datawizard/sync_and_notify_data_implicit.c

@@ -49,8 +49,10 @@ void cuda_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args);
 #endif
 
 #ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
 void opencl_codelet_incA(void *descr[], __attribute__ ((unused)) void *_args);
 void opencl_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args);
+struct starpu_opencl_codelet opencl_code;
 #endif
 
 #define VECTORSIZE	16
@@ -120,7 +122,7 @@ int main(int argc, char **argv)
 #endif
 
 #ifdef STARPU_USE_OPENCL
-        _starpu_opencl_compile_source_to_opencl("tests/datawizard/sync_and_notify_data_opencl_codelet.cl");
+        starpu_opencl_load_opencl_from_file("tests/datawizard/sync_and_notify_data_opencl_codelet.cl", &opencl_code);
 #endif
 
         starpu_vector_data_register(&v_handle, 0, (uintptr_t)v, VECTORSIZE, sizeof(unsigned));

tests/datawizard/sync_and_notify_data_opencl.c

@@ -18,6 +18,8 @@
 #include <starpu_opencl.h>
 #include <CL/cl.h>
 
+extern struct starpu_opencl_codelet opencl_code;
+
 void opencl_codelet_incA(void *descr[], __attribute__ ((unused)) void *_args)
 {
         unsigned *val = (unsigned *)STARPU_GET_VECTOR_PTR(descr[0]);
@@ -28,7 +30,7 @@ void opencl_codelet_incA(void *descr[], __attribute__ ((unused)) void *_args)
 	id = starpu_worker_get_id();
 	devid = starpu_worker_get_devid(id);
 
-	err = starpu_opencl_load_kernel(&kernel, &queue, "tests/datawizard/sync_and_notify_data_opencl_codelet.cl", "incA", devid);
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "incA", devid);
 	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -44,7 +46,7 @@ void opencl_codelet_incA(void *descr[], __attribute__ ((unused)) void *_args)
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }
 
 void opencl_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args)
@@ -57,7 +59,7 @@ void opencl_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args)
 	id = starpu_worker_get_id();
 	devid = starpu_worker_get_devid(id);
 
-	err = starpu_opencl_load_kernel(&kernel, &queue, "tests/datawizard/sync_and_notify_data_opencl_codelet.cl", "incC", devid);
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_code, "incC", devid);
 	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
 
 	err = 0;
@@ -73,5 +75,5 @@ void opencl_codelet_incC(void *descr[], __attribute__ ((unused)) void *_args)
 
 	clFinish(queue);
 
-	starpu_opencl_release(kernel);
+	starpu_opencl_release_kernel(kernel);
 }