[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