[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [32194] branches/particles-2010/source/ blender/nodes/intern/simulation: Better solution for accessing socket buffers from kernels and implicit conversions of input data .
Lukas Toenne
lukas.toenne at googlemail.com
Wed Sep 29 18:32:38 CEST 2010
Revision: 32194
http://projects.blender.org/plugins/scmsvn/viewcvs.php?view=rev&root=bf-blender&revision=32194
Author: lukastoenne
Date: 2010-09-29 18:32:38 +0200 (Wed, 29 Sep 2010)
Log Message:
-----------
Better solution for accessing socket buffers from kernels and implicit conversions of input data. There is now a set of macros for getting data from input buffers and writing to output buffers. Similar versions exist for both the C kernel functions and OpenCL code.
Modified Paths:
--------------
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_math.c
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-29 16:18:39 UTC (rev 32193)
+++ branches/particles-2010/source/blender/nodes/intern/simulation/SIM_util.c 2010-09-29 16:32:38 UTC (rev 32194)
@@ -71,227 +71,12 @@
);
/* 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)(x > 0.0f)\n"
-"#define CONVERT_BOOL_INT(x) (char)(x > 0)\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_macros = "";
-const char *cl_socket_funcs = STRINGIFY(
-size_t get_socket_id(int size)
-{
- return (get_global_id(0) % (size_t)size);
-}
-);
+const char *cl_socket_funcs = "";
#endif
-static void read_float_from_float(int index, void *data_arg, void *size_arg, void *result)
-{
- *(float*)result = *(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg));
-}
-
-static void read_float_from_int(int index, void *data_arg, void *size_arg, void *result)
-{
- *(float*)result = (float)(*(int*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg)));
-}
-
-static void read_float_from_bool(int index, void *data_arg, void *size_arg, void *result)
-{
- *(float*)result = (float)(*(char*)(*(char**)data_arg + SIM_ITEMSIZE_BOOL * (index % *(int*)size_arg)));
-}
-
-static void read_int_from_float(int index, void *data_arg, void *size_arg, void *result)
-{
- *(int*)result = (int)(*(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg)));
-}
-
-static void read_int_from_int(int index, void *data_arg, void *size_arg, void *result)
-{
- *(int*)result = *(float*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg));
-}
-
-static void read_int_from_bool(int index, void *data_arg, void *size_arg, void *result)
-{
- *(int*)result = (int)(*(*(char**)data_arg + SIM_ITEMSIZE_BOOL * (index % *(int*)size_arg)));
-}
-
-static void read_bool_from_float(int index, void *data_arg, void *size_arg, void *result)
-{
- *(char*)result = (*(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg)) > 0.0f);
-}
-
-static void read_bool_from_int(int index, void *data_arg, void *size_arg, void *result)
-{
- *(char*)result = (*(int*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg)) > 0);
-}
-
-static void read_bool_from_bool(int index, void *data_arg, void *size_arg, void *result)
-{
- *(char*)result = *(char*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg));
-}
-
-static void read_vector_from_vector(int index, void *data_arg, void *size_arg, void *result)
-{
- *(float**)result = (float*)(*(char**)data_arg + SIM_ITEMSIZE_VECTOR * (index % *(int*)size_arg));
-}
-
-static void read_rgba_from_rgba(int index, void *data_arg, void *size_arg, void *result)
-{
- *(float**)result = (float*)(*(char**)data_arg + SIM_ITEMSIZE_RGBA * (index % *(int*)size_arg));
-}
-
-static void read_string_from_string(int index, void *data_arg, void *size_arg, void *result)
-{
- *(char**)result = (char*)(*(char**)data_arg + SIM_ITEMSIZE_STRING * (index % *(int*)size_arg));
-}
-
-static SimBufferReadFunction find_read_function(int fromtype, int totype)
-{
- switch (fromtype) {
- case SOCK_FLOAT:
- switch (totype) {
- case SOCK_FLOAT: return read_float_from_float; break;
- case SOCK_INT: return read_int_from_float; break;
- case SOCK_BOOL: return read_bool_from_float; break;
- }
- break;
- case SOCK_INT:
- switch (totype) {
- case SOCK_FLOAT: return read_float_from_int; break;
- case SOCK_INT: return read_int_from_int; break;
- case SOCK_BOOL: return read_bool_from_int; break;
- }
- break;
- case SOCK_BOOL:
- switch (totype) {
- case SOCK_FLOAT: return read_float_from_bool; break;
- case SOCK_INT: return read_int_from_bool; break;
- case SOCK_BOOL: return read_bool_from_bool; break;
- }
- break;
- case SOCK_VECTOR:
- if (totype == SOCK_VECTOR) return read_vector_from_vector;
- break;
- case SOCK_RGBA:
- if (totype == SOCK_RGBA) return read_rgba_from_rgba;
- break;
- case SOCK_STRING:
- if (totype == SOCK_STRING) return read_string_from_string;
- break;
- }
- return NULL;
-}
-
-static void write_float_to_float(int index, void *data_arg, void *size_arg, void *value)
-{
- *(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg)) = *(float*)value;
-}
-
-static void write_float_to_int(int index, void *data_arg, void *size_arg, void *value)
-{
- *(int*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg)) = (int)(*(float*)value);
-}
-
-static void write_float_to_bool(int index, void *data_arg, void *size_arg, void *value)
-{
- *(char*)(*(char**)data_arg + SIM_ITEMSIZE_BOOL * (index % *(int*)size_arg)) = (*(float*)value > 0.0f);
-}
-
-static void write_int_to_float(int index, void *data_arg, void *size_arg, void *value)
-{
- *(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg)) = (float)(*(int*)value);
-}
-
-static void write_int_to_int(int index, void *data_arg, void *size_arg, void *value)
-{
- *(int*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg)) = *(int*)value;
-}
-
-static void write_int_to_bool(int index, void *data_arg, void *size_arg, void *value)
-{
- *(char*)(*(char**)data_arg + SIM_ITEMSIZE_BOOL * (index % *(int*)size_arg)) = (*(int*)value > 0);
-}
-
-static void write_bool_to_float(int index, void *data_arg, void *size_arg, void *value)
-{
- *(float*)(*(char**)data_arg + SIM_ITEMSIZE_FLOAT * (index % *(int*)size_arg)) = (float)(*(char*)value);
-}
-
-static void write_bool_to_int(int index, void *data_arg, void *size_arg, void *value)
-{
- *(int*)(*(char**)data_arg + SIM_ITEMSIZE_INT * (index % *(int*)size_arg)) = (int)(*(char*)value);
-}
-
-static void write_bool_to_bool(int index, void *data_arg, void *size_arg, void *value)
-{
- *(char*)(*(char**)data_arg + SIM_ITEMSIZE_BOOL * (index % *(int*)size_arg)) = *(char*)value;
-}
-
-static void write_vector_to_vector(int index, void *data_arg, void *size_arg, void *value)
-{
- copy_v3_v3((float*)(*(char**)data_arg + SIM_ITEMSIZE_VECTOR * (index % *(int*)size_arg)), *(float**)value);
-}
-
-static void write_rgba_to_rgba(int index, void *data_arg, void *size_arg, void *value)
-{
- float *res = (float*)(*(char**)data_arg + SIM_ITEMSIZE_RGBA * (index % *(int*)size_arg));
- res[0] = (*(float**)value)[0];
- res[1] = (*(float**)value)[1];
- res[2] = (*(float**)value)[2];
- res[3] = (*(float**)value)[3];
-}
-
-static void write_string_to_string(int index, void *data_arg, void *size_arg, void *value)
-{
- memcpy((char*)(*(char**)data_arg + SIM_ITEMSIZE_STRING * (index % *(int*)size_arg)), *(char**)value, SIM_ITEMSIZE_STRING);
-}
-
-static SimBufferWriteFunction find_write_function(int fromtype, int totype)
-{
- switch (fromtype) {
- case SOCK_FLOAT:
- switch (totype) {
- case SOCK_FLOAT: return write_float_to_float; break;
- case SOCK_INT: return write_float_to_int; break;
- case SOCK_BOOL: return write_float_to_bool; break;
- }
- break;
- case SOCK_INT:
- switch (totype) {
- case SOCK_FLOAT: return write_int_to_float; break;
- case SOCK_INT: return write_int_to_int; break;
- case SOCK_BOOL: return write_int_to_bool; break;
- }
- break;
- case SOCK_BOOL:
- switch (totype) {
- case SOCK_FLOAT: return write_bool_to_float; break;
- case SOCK_INT: return write_bool_to_int; break;
- case SOCK_BOOL: return write_bool_to_bool; break;
- }
- break;
- case SOCK_VECTOR:
- if (totype == SOCK_VECTOR) return write_vector_to_vector;
- break;
- case SOCK_RGBA:
- if (totype == SOCK_RGBA) return write_rgba_to_rgba;
- break;
- case SOCK_STRING:
- if (totype == SOCK_STRING) return write_string_to_string;
- break;
- }
- return NULL;
-}
-
size_t sim_get_data_size(int datatype)
{
#ifdef WITH_OPENCL
@@ -1048,6 +833,61 @@
}
}
+void sim_kernel_set_node(SimExecData *execdata, SimKernel *kernel, SimNodeStack *node, int num_args)
+{
+ if (execdata->error)
+ return;
+
+ kernel->node = node;
+
+#ifdef WITH_OPENCL
+ if (BKE_opencl_is_active()) {
+ int i, arg_index;
+ cl_int res;
+ /* Append socket buffer arguments after custom arguments.
+ * This includes the buffer data pointer as well as the context size (number of elements).
+ */
+ arg_index = num_args;
+ for (i=0; i < node->totin; ++i) {
+ /* note: the buffer data is casted to actual data type pointer in OpenCL code (float*, int*, etc.),
+ * but can be handled as void* here.
+ */
+ res = clSetKernelArg(kernel->impl_cl, arg_index++, sizeof(cl_mem), &node->instack[i]->buffer.data_cl);
+ if (res != CL_SUCCESS) {
+ execdata->error = 1;
+ sprintf(execdata->error_string, "Error setting OpenCL kernel input socket buffer argument: %s", BKE_opencl_message(res));
+ return;
+ }
+
+ res = clSetKernelArg(kernel->impl_cl, arg_index++, sizeof(int), &node->instack[i]->context.size);
+ if (res != CL_SUCCESS) {
+ execdata->error = 1;
+ sprintf(execdata->error_string, "Error setting OpenCL kernel input socket size argument: %s", BKE_opencl_message(res));
+ return;
+ }
+ }
+ for (i=0; i < node->totout; ++i) {
+ res = clSetKernelArg(kernel->impl_cl, arg_index++, sizeof(cl_mem), &node->outstack[i].buffer.data_cl);
+ if (res != CL_SUCCESS) {
+ execdata->error = 1;
+ sprintf(execdata->error_string, "Error setting OpenCL kernel output socket buffer argument: %s", BKE_opencl_message(res));
+ return;
+ }
+
+ res = clSetKernelArg(kernel->impl_cl, arg_index++, sizeof(int), &node->outstack[i].context.size);
+ if (res != CL_SUCCESS) {
+ execdata->error = 1;
+ sprintf(execdata->error_string, "Error setting OpenCL kernel output socket size argument: %s", BKE_opencl_message(res));
+ return;
+ }
+ }
+ }
+ else
+#endif
+ {
+ }
+}
+
void sim_set_kernel_arg(SimExecData *execdata, SimKernel *kernel, int arg_index, size_t arg_size, const void *arg_value)
{
if (execdata->error)
@@ -1073,6 +913,7 @@
}
}
+#if 0
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list