[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [31963] branches/particles-2010/source/ blender: Added a wrapper for OpenCL kernels and implemented the "Add" node for testing.

Lukas Toenne lukas.toenne at googlemail.com
Thu Sep 16 11:53:04 CEST 2010


Revision: 31963
          http://projects.blender.org/plugins/scmsvn/viewcvs.php?view=rev&root=bf-blender&revision=31963
Author:   lukastoenne
Date:     2010-09-16 11:53:04 +0200 (Thu, 16 Sep 2010)

Log Message:
-----------
Added a wrapper for OpenCL kernels and implemented the "Add" node for testing.

Modified Paths:
--------------
    branches/particles-2010/source/blender/blenkernel/BKE_node.h
    branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.c
    branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.h
    branches/particles-2010/source/blender/nodes/intern/simulation/node_tree_simulation.c
    branches/particles-2010/source/blender/nodes/intern/simulation/nodes/SIM_get_data.c
    branches/particles-2010/source/blender/nodes/intern/simulation/nodes/SIM_if.c
    branches/particles-2010/source/blender/nodes/intern/simulation/nodes/SIM_math.c
    branches/particles-2010/source/blender/nodes/intern/simulation/nodes/SIM_program.c
    branches/particles-2010/source/blender/nodes/intern/simulation/nodes/SIM_set_data.c

Modified: branches/particles-2010/source/blender/blenkernel/BKE_node.h
===================================================================
--- branches/particles-2010/source/blender/blenkernel/BKE_node.h	2010-09-16 07:26:06 UTC (rev 31962)
+++ branches/particles-2010/source/blender/blenkernel/BKE_node.h	2010-09-16 09:53:04 UTC (rev 31963)
@@ -122,7 +122,7 @@
 	void (*freesocketstoragefunc)(struct bNodeSocket *sock);
 
 	/* simulation nodes */
-	const char *cl_source;
+	char *(*generate_source)(struct SimNodeStack *node);
 	void (*enqueue)(struct SimExecData *execdata, struct SimNodeStack *node, struct SimDataContext *self);
 	void (*enqueue_op)(struct SimExecData *execdata, struct SimNodeStack *node, struct SimDataContext *self, int execlevel, int *pushop);
 } bNodeType;

Modified: branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.c
===================================================================
--- branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.c	2010-09-16 07:26:06 UTC (rev 31962)
+++ branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.c	2010-09-16 09:53:04 UTC (rev 31963)
@@ -44,27 +44,50 @@
 #include "RNA_access.h"
 
 #ifdef WITH_OPENCL
+
 const char *cl_string_funcs = STRINGIFY(
-size_t strlen(const char *str)
-{
-	char *end= str;
-	while (*end != '\0')
-		++end;
-	return (end - str);
-}
+//size_t strlen(const char *str)
+//{
+//	const char *end= str;
+//	while (*end != '\0')
+//		++end;
+//	return (end - str);
+//}
 
-char *strcpy(char *destination, const char *source)
+//char *strcpy(char *destination, const char *source)
+//{
+//	while (*source != '\0') {
+//		*destination = *source;
+//		++source;
+//		++destination;
+//	}
+//	*destination = '\0';
+//	return destination;
+//}
+);
+
+/* note: can't use stringify here because of precompiler definitions */
+const char *cl_socket_macros = 
+"#define CONVERT_FLOAT_FLOAT(x)		(x)\n"
+"#define CONVERT_FLOAT_INT(x)		(float)(x)\n"
+"#define CONVERT_FLOAT_BOOL(x)		(float)(x)\n"
+"#define CONVERT_INT_FLOAT(x)		(int)(x)\n"
+"#define CONVERT_INT_INT(x)			(x)\n"
+"#define CONVERT_INT_BOOL(x)		(int)(x)\n"
+"#define CONVERT_BOOL_FLOAT(x)		(char)((int)(x) % 2)\n"
+"#define CONVERT_BOOL_INT(x)		(char)((x) % 2)\n"
+"#define CONVERT_BOOL_BOOL(x)		(x)\n"
+"#define CONVERT_VECTOR_VECTOR(x)	(x)\n"
+"#define CONVERT_RGBA_RGBA(x)		(x)\n"
+"#define CONVERT_STRING_STRING(x)	(x)\n";
+
+const char *cl_socket_funcs = STRINGIFY(
+size_t get_socket_id(int size)
 {
-	while (*source != '\0') {
-		*destination = *source;
-		++source;
-		++destination;
-	}
-	*destination = '\0';
-	return destination;
+	return (get_global_id(0) % (size_t)size);
 }
+);
 
-);
 #endif
 
 size_t sim_get_data_size(int datatype)
@@ -117,6 +140,52 @@
 	#endif
 }
 
+const char *sim_get_typename(int datatype)
+{
+	switch (datatype) {
+	case SOCK_FLOAT:
+		return "FLOAT";
+	case SOCK_INT:
+		return "INT";
+	case SOCK_BOOL:
+		return "BOOL";
+	case SOCK_VECTOR:
+		return "VECTOR";
+	case SOCK_RGBA:
+		return "RGBA";
+	case SOCK_STRING:
+		return "STRING";
+	case SOCK_OP:
+		return "OP";
+	default:
+		return "";
+	}
+}
+
+#ifdef WITH_OPENCL
+const char *sim_get_opencl_type(int datatype)
+{
+	switch (datatype) {
+	case SOCK_FLOAT:
+		return "float";
+	case SOCK_INT:
+		return "int";
+	case SOCK_BOOL:
+		return "char";
+	case SOCK_VECTOR:
+		return "float4";
+	case SOCK_RGBA:
+		return "float4";
+	case SOCK_STRING:
+		return "char*";
+	case SOCK_OP:
+		return "int";
+	default:
+		return "int";
+	}
+}
+#endif
+
 static void context_update_size(SimDataContext *ctx)
 {
 	if (ctx->prop && RNA_property_type(ctx->prop) == PROP_COLLECTION) {
@@ -578,7 +647,161 @@
 }
 #endif
 
+SimKernel sim_create_kernel(SimExecData *execdata, KernelFunction func, const char *cl_funcname)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res;
+		SimKernel kernel;
+		
+		kernel.impl_cl = clCreateKernel(*execdata->program, cl_funcname, &res);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error creating OpenCL kernel: %s", BKE_opencl_message(res));
+		}
+		
+		return kernel;
+	}
+	else {
+#endif
+		SimKernel kernel;
+		int i;
+		
+		kernel.impl = func;
+		for (i=0; i < SIM_MAXKERNELARGS; ++i) {
+			kernel.args[i] = NULL;
+		}
+		return kernel;
+#ifdef WITH_OPENCL
+	}
+#endif
+}
 
+void sim_release_kernel(SimExecData *execdata, SimKernel *kernel)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res = clReleaseKernel(kernel->impl_cl);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error releasing OpenCL kernel: %s", BKE_opencl_message(res));
+		}
+	}
+	else {
+#endif
+		int i;
+		
+		for (i=0; i < SIM_MAXKERNELARGS; ++i) {
+			if (kernel->args[i]) {
+				MEM_freeN(kernel->args[i]);
+				kernel->args[i] = NULL;
+			}
+		}
+#ifdef WITH_OPENCL
+	}
+#endif
+}
+
+void sim_set_kernel_arg(SimExecData *execdata, SimKernel *kernel, int arg_index, size_t arg_size, const void *arg_value)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res;
+		res = clSetKernelArg(kernel->impl_cl, arg_index, arg_size, arg_value);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error setting OpenCL kernel argument: %s", BKE_opencl_message(res));
+		}
+	}
+	else {
+#endif
+		assert( arg_index < SIM_MAXKERNELARGS );
+		if (kernel->args[arg_index])
+			MEM_freeN(kernel->args[arg_index]);
+		kernel->args[arg_index] = MEM_callocN(arg_size, "kernel argument");
+		memcpy(kernel->args[arg_index], arg_value, arg_size);
+#ifdef WITH_OPENCL
+	}
+#endif
+}
+
+void sim_set_kernel_socket_data_arg(SimExecData *execdata, SimKernel *kernel, int arg_index, SimSocketStack *socket)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res;
+		res = clSetKernelArg(kernel->impl_cl, arg_index, sizeof(cl_mem), &socket->buffer.data_cl);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error setting OpenCL kernel socket data argument: %s", BKE_opencl_message(res));
+		}
+	}
+	else {
+#endif
+		assert( arg_index < SIM_MAXKERNELARGS );
+		if (kernel->args[arg_index])
+			MEM_freeN(kernel->args[arg_index]);
+		kernel->args[arg_index] = MEM_callocN(sizeof(void*), "kernel socket data argument");
+		memcpy(kernel->args[arg_index], &socket->buffer.data, sizeof(void*));
+#ifdef WITH_OPENCL
+	}
+#endif
+}
+
+void sim_set_kernel_socket_size_arg(SimExecData *execdata, SimKernel *kernel, int arg_index, SimSocketStack *socket)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res;
+		res = clSetKernelArg(kernel->impl_cl, arg_index, sizeof(int), &socket->context.size);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error setting OpenCL kernel socket size argument: %s", BKE_opencl_message(res));
+		}
+	}
+	else {
+#endif
+		assert( arg_index < SIM_MAXKERNELARGS );
+		if (kernel->args[arg_index])
+			MEM_freeN(kernel->args[arg_index]);
+		kernel->args[arg_index] = MEM_callocN(sizeof(int), "kernel socket size argument");
+		memcpy(kernel->args[arg_index], &socket->context.size, sizeof(int));
+#ifdef WITH_OPENCL
+	}
+#endif
+}
+
+void sim_enqueue_kernel(struct SimExecData *execdata, SimKernel *kernel, size_t total, int num_wait_events, SimEvent *wait_events, SimEvent *event)
+{
+#ifdef WITH_OPENCL
+	if (BKE_opencl_is_active()) {
+		cl_int res;
+		cl_event *wait_events_cl;
+		
+		if (num_wait_events > 0) {
+			int i;
+			wait_events_cl = MEM_callocN(num_wait_events * sizeof(cl_event), "wait_events_cl");
+			for (i=0; i < num_wait_events; ++i)
+				wait_events_cl[i] = wait_events[i].impl_cl;
+		}
+		else
+			wait_events_cl = NULL;
+		res = clEnqueueNDRangeKernel(execdata->queue, kernel->impl_cl, 1, NULL, &total, NULL, num_wait_events, wait_events_cl, &event->impl_cl);
+		MEM_freeN(wait_events_cl);
+		if (res != CL_SUCCESS) {
+			execdata->error = 1;
+			sprintf(execdata->error_string, "Error enqueueing OpenCL kernel: %s", BKE_opencl_message(res));
+		}
+	}
+	else {
+#endif
+		/* TODO */
+#ifdef WITH_OPENCL
+	}
+#endif
+}
+
+
 int sim_node_prepare_output(struct SimExecData *execdata, struct SimSocketStack *out, struct SimDataContext *ctx)
 {
 	out->context = *ctx;
@@ -612,11 +835,11 @@
 					res = *ctx;
 				}
 				else {
-					res.scene = NULL;
-					RNA_pointer_create(NULL, NULL, NULL, &res.ptr);
-					res.prop = NULL;
-					res.key = 0;
-					res.size = 0;
+					r_max->scene = NULL;
+					RNA_pointer_create(NULL, NULL, NULL, &r_max->ptr);
+					r_max->prop = NULL;
+					r_max->key = 0;
+					r_max->size = 0;
 					return 0;
 				}
 			}
@@ -632,6 +855,7 @@
 			}
 		}
 	}
+	*r_max = res;
 	return 1;
 }
 

Modified: branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.h
===================================================================
--- branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.h	2010-09-16 07:26:06 UTC (rev 31962)
+++ branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.h	2010-09-16 09:53:04 UTC (rev 31963)
@@ -62,8 +62,12 @@
 
 #define SIM_STRINGLENGTH		128
 #define SIM_MAXDATASIZE			128		/* maximum size of any data element type */
+#define SIM_MAXKERNELARGS		32
 
+/* common OpenCL helper functions */
 extern const char *cl_string_funcs;
+extern const char *cl_socket_funcs;
+extern const char *cl_socket_macros;
 
 typedef struct bNodeSocketIsland {
 	struct bNodeSocketIsland *next, *prev;
@@ -149,6 +153,16 @@
 	#endif
 } SimEvent;
 
+typedef void (*KernelFunction)(void **args);
+
+typedef struct SimKernel {
+	#ifdef WITH_OPENCL
+	cl_kernel impl_cl;
+	#endif
+	KernelFunction impl;
+	void *args[SIM_MAXKERNELARGS];
+} SimKernel;
+
 struct SimNodeStack;
 typedef struct SimSocketStack {
 	bNodeSocket *base;
@@ -170,8 +184,9 @@
 typedef struct SimNodeStack {
 	SimNodeState state;
 	short flags;
-	
 	bNode *base;
+	int id;
+
 	struct SimSocketStack **instack;
 	struct SimSocketStack *outstack;
 	int totin, totout;
@@ -210,6 +225,10 @@
 } SimExecData;
 
 size_t sim_get_data_size(int datatype);
+const char *sim_get_typename(int datatype);
+#ifdef WITH_OPENCL
+const char *sim_get_opencl_type(int datatype);
+#endif
 
 void sim_context_create(struct Scene *scene, struct PointerRNA *ptr, struct PropertyRNA *prop, int key, struct SimDataContext *r_ctx);
 struct StructRNA *sim_context_type(struct SimDataContext *ctx);
@@ -232,6 +251,13 @@
 //SimEvent sim_create_user_event(struct SimExecData *execdata);
 //void sim_set_user_event_status(SimEvent event, SimEventStatus status);
 

@@ Diff output truncated at 10240 characters. @@




More information about the Bf-blender-cvs mailing list