[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [31882] branches/particles-2010: OpenCL program is now compiled when the tree is executed and only when the tree has actually changed ( flag is currently disabled because program does not depend on tree state anyway ).
Lukas Toenne
lukas.toenne at googlemail.com
Sat Sep 11 16:43:28 CEST 2010
Revision: 31882
http://projects.blender.org/plugins/scmsvn/viewcvs.php?view=rev&root=bf-blender&revision=31882
Author: lukastoenne
Date: 2010-09-11 16:43:27 +0200 (Sat, 11 Sep 2010)
Log Message:
-----------
OpenCL program is now compiled when the tree is executed and only when the tree has actually changed (flag is currently disabled because program does not depend on tree state anyway). The compilation can also be forced by the user explicitly (button in node editor header). This avoids delays due to OpenCL compilation at startup time and gives users more control, while still doing automatic compiles when necessary.
Modified Paths:
--------------
branches/particles-2010/release/scripts/ui/space_node.py
branches/particles-2010/source/blender/blenkernel/BKE_node.h
branches/particles-2010/source/blender/blenkernel/intern/node.c
branches/particles-2010/source/blender/blenloader/intern/readfile.c
branches/particles-2010/source/blender/editors/space_node/node_edit.c
branches/particles-2010/source/blender/editors/space_node/node_intern.h
branches/particles-2010/source/blender/editors/space_node/node_ops.c
branches/particles-2010/source/blender/makesdna/DNA_node_types.h
branches/particles-2010/source/blender/makesrna/intern/rna_space.c
branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_get_data.c
branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_math.c
branches/particles-2010/source/blender/nodes/intern/SIM_util.c
branches/particles-2010/source/blender/nodes/intern/SIM_util.h
branches/particles-2010/source/blender/nodes/intern/node_tree_simulation.c
Modified: branches/particles-2010/release/scripts/ui/space_node.py
===================================================================
--- branches/particles-2010/release/scripts/ui/space_node.py 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/release/scripts/ui/space_node.py 2010-09-11 14:43:27 UTC (rev 31882)
@@ -69,6 +69,12 @@
layout.prop(scene.render, "use_free_unused_nodes", text="Free Unused")
layout.prop(snode, "show_backdrop")
+ elif snode.tree_type == 'SIMULATION':
+ scene = snode.id
+
+ props = layout.operator("node.compile_opencl_program", text="Compile Program")
+ props.force = True
+
layout.separator()
layout.template_running_jobs()
Modified: branches/particles-2010/source/blender/blenkernel/BKE_node.h
===================================================================
--- branches/particles-2010/source/blender/blenkernel/BKE_node.h 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/blenkernel/BKE_node.h 2010-09-11 14:43:27 UTC (rev 31882)
@@ -564,6 +564,8 @@
#define SIM_NODE_DEBUGPRINT 801
/* API */
+void ntreeSimulationCompileProgram(struct bNodeTree *ntree, int force);
+void ntreeSimulationReleaseProgram(struct bNodeTree *ntree);
void ntreeSimulationExecTree(struct bNodeTree *ntree, struct Scene *scene, struct PointerRNA *self, float frame, float dframe, float time, float dtime);
struct StructRNA *simnode_getdata_get_type(struct bNode *node);
Modified: branches/particles-2010/source/blender/blenkernel/intern/node.c
===================================================================
--- branches/particles-2010/source/blender/blenkernel/intern/node.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/blenkernel/intern/node.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -1257,8 +1257,10 @@
MEM_freeN(ntree->owntype);
}
- if (ntree->islands) {
- MEM_freeN(ntree->islands);
+ if (ntree->type == NTREE_SIMULATION) {
+ if (ntree->islands)
+ MEM_freeN(ntree->islands);
+ ntreeSimulationReleaseProgram(ntree);
}
}
Modified: branches/particles-2010/source/blender/blenloader/intern/readfile.c
===================================================================
--- branches/particles-2010/source/blender/blenloader/intern/readfile.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/blenloader/intern/readfile.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -2137,6 +2137,7 @@
ntree->init= 0; /* to set callbacks and force setting types */
ntree->owntype= NULL;
ntree->progress= NULL;
+ ntree->program= NULL;
ntree->adt= newdataadr(fd, ntree->adt);
direct_link_animdata(fd, ntree->adt);
Modified: branches/particles-2010/source/blender/editors/space_node/node_edit.c
===================================================================
--- branches/particles-2010/source/blender/editors/space_node/node_edit.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/editors/space_node/node_edit.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -2516,3 +2516,44 @@
but= uiDefSearchBut(block, search, 0, ICON_VIEWZOOM, sizeof(search), 0, 0, UI_UNIT_X*6, UI_UNIT_Y, 0, 0, "");
uiButSetSearchFunc(but, simnode_setdata_add_rna_socket_search_cb, node, simnode_setdata_add_rna_socket_exec_cb, NULL);
}
+
+/* ****************** Compile OpenCL Program Node Operator ******************* */
+
+static int node_compile_opencl_program_exec(bContext *C, wmOperator *op)
+{
+ SpaceNode *snode= CTX_wm_space_node(C);
+ int force = 0;
+
+ /* check input variables */
+ if (RNA_property_is_set(op->ptr, "force"))
+ {
+ force = RNA_boolean_get(op->ptr, "force");
+ }
+
+ /* TODO necessary? */
+// ED_preview_kill_jobs(C);
+
+ ntreeSimulationCompileProgram(snode->nodetree, force);
+
+ snode_notify(C, snode);
+
+ return OPERATOR_FINISHED;
+}
+
+void NODE_OT_compile_opencl_program(wmOperatorType *ot)
+{
+ /* identifiers */
+ ot->name= "Compile node tree program";
+ ot->description= "Compile the OpenCL program for the execution of a simulation node tree";
+ ot->idname= "NODE_OT_compile_opencl_program";
+
+ /* callbacks */
+ ot->exec= node_compile_opencl_program_exec;
+ ot->poll= ED_operator_node_active;
+
+ /* flags */
+ ot->flag= OPTYPE_REGISTER;
+
+ RNA_def_boolean(ot->srna, "force", 0, "force", "Force the compilation, whether the tree has changed or not.");
+}
+
Modified: branches/particles-2010/source/blender/editors/space_node/node_intern.h
===================================================================
--- branches/particles-2010/source/blender/editors/space_node/node_intern.h 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/editors/space_node/node_intern.h 2010-09-11 14:43:27 UTC (rev 31882)
@@ -122,6 +122,8 @@
void NODE_OT_add_file(struct wmOperatorType *ot);
+void NODE_OT_compile_opencl_program(struct wmOperatorType *ot);
+
// XXXXXX
// XXX from BSE_node.h
Modified: branches/particles-2010/source/blender/editors/space_node/node_ops.c
===================================================================
--- branches/particles-2010/source/blender/editors/space_node/node_ops.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/editors/space_node/node_ops.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -83,6 +83,8 @@
WM_operatortype_append(NODE_OT_backimage_zoom);
WM_operatortype_append(NODE_OT_add_file);
+
+ WM_operatortype_append(NODE_OT_compile_opencl_program);
}
void ED_operatormacros_node(void)
Modified: branches/particles-2010/source/blender/makesdna/DNA_node_types.h
===================================================================
--- branches/particles-2010/source/blender/makesdna/DNA_node_types.h 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/makesdna/DNA_node_types.h 2010-09-11 14:43:27 UTC (rev 31882)
@@ -212,7 +212,10 @@
struct bNodeType *owntype; /* for groups or dynamic trees, no read/write */
+ /* TODO move this to a separate storage struct to keep the tree free of type-dependent stuff? */
+ /* simulation trees */
struct bNodeSocketIsland *islands; /* used during calculation of socket type islands */
+ void *program; /* OpenCL compiled program (in case it's available and activated) */
/* execution data */
/* XXX It would be preferable to completely move this data out of the underlying node tree,
@@ -239,12 +242,13 @@
#define NUM_NTREE_TYPES 4
/* ntree->init, flag */
-#define NTREE_TYPE_INIT 1
+#define NTREE_TYPE_INIT 1
/* ntree->flag */
#define NTREE_DS_EXPAND 1 /* for animation editors */
#define NTREE_UPDATE 2 /* tagged for update */
#define NTREE_TAGGED 4 /* generic tag for recursive functions */
+#define NTREE_NEED_RECOMPILE 8 /* simulation tree has program changes that require recompile */
/* data structs, for node->storage */
Modified: branches/particles-2010/source/blender/makesrna/intern/rna_space.c
===================================================================
--- branches/particles-2010/source/blender/makesrna/intern/rna_space.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/makesrna/intern/rna_space.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -2176,7 +2176,7 @@
{NTREE_SHADER, "MATERIAL", ICON_MATERIAL, "Material", "Material nodes"},
{NTREE_TEXTURE, "TEXTURE", ICON_TEXTURE, "Texture", "Texture nodes"},
{NTREE_COMPOSIT, "COMPOSITING", ICON_RENDERLAYERS, "Compositing", "Compositing nodes"},
- {NTREE_SIMULATION, "PARTICLES", ICON_PHYSICS, "Simulation", "Simulation nodes"},
+ {NTREE_SIMULATION, "SIMULATION", ICON_PHYSICS, "Simulation", "Simulation nodes"},
{0, NULL, 0, NULL, NULL}};
static EnumPropertyItem texture_type_items[] = {
Modified: branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_get_data.c
===================================================================
--- branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_get_data.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_get_data.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -176,7 +176,7 @@
{
size_t datasize = sim_get_data_size(out->type);
- if (!sim_node_prepare_output(execdata, out, ctx, 0) || execdata->error)
+ if (!sim_node_prepare_output(execdata, out, ctx) || execdata->error)
return;
if (sim_context_is_collection(ctx)) {
@@ -292,7 +292,7 @@
SimDataContext ctx;
int resolved;
- if (sim_node_prepare_output(execdata, &node->outstack[0], &node->instack[0]->context, 0)) {
+ if (sim_node_prepare_output(execdata, &node->outstack[0], &node->instack[0]->context)) {
if (execdata->error) return;
/* copy the path socket to the output */
sim_enqueue_copy_buffer(execdata, &node->instack[0]->buffer, &node->outstack[0].buffer, 0, 0, 1, node->totin, node->event_wait_list, &node->outstack[0].event);
Modified: branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_math.c
===================================================================
--- branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_math.c 2010-09-11 11:16:55 UTC (rev 31881)
+++ branches/particles-2010/source/blender/nodes/intern/SIM_nodes/SIM_math.c 2010-09-11 14:43:27 UTC (rev 31882)
@@ -78,12 +78,36 @@
static const char cl_source_add[]= STRINGIFY(
-__kernel void node_add(__global float *res, __global float *a, __global float *b)
+__kernel void node_add_float(__global float *res, __global float *a, __global float *b)
{
unsigned int i = get_global_id(0);
res[i] = a[i] + b[i];
}
+__kernel void node_add_int(__global int *res, __global int *a, __global int *b)
+{
+ unsigned int i = get_global_id(0);
+ res[i] = a[i] + b[i];
+}
+
+__kernel void node_add_bool(__global char *res, __global char *a, __global char *b)
+{
+ unsigned int i = get_global_id(0);
+ res[i] = a[i] || b[i];
+}
+
+__kernel void node_add_vector(__global float4 *res, __global float4 *a, __global float4 *b)
+{
+ unsigned int i = get_global_id(0);
+ res[i] = a[i] + b[i];
+}
+
+__kernel void node_add_rgba(__global float4 *res, __global float4 *a, __global float4 *b)
+{
+ unsigned int i = get_global_id(0);
+ res[i] = a[i] + b[i];
+}
+
);
static void cl_enqueue_add(SimExecData *execdata, SimNodeStack *node, SimDataContext *self)
@@ -93,10 +117,16 @@
int i;
/* find the input context */
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list