diff options
| author | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-06-12 16:09:33 +0300 |
|---|---|---|
| committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-06-12 16:12:36 +0300 |
| commit | 0f4f4d87542b96942234e10421823dcc6766331c (patch) | |
| tree | fd545e6a4ed9ad0652e7e7e18fb76a9e1324e748 | |
| parent | f52dc2f371923c22a974df7105245f7e0b8148ee (diff) | |
| parent | d8957e4ccec4cda1cd72e94045efba6ab2ceb6fd (diff) | |
Merge branch 'master' into blender2.8
114 files changed, 1533 insertions, 481 deletions
diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py index 9e22ce331d0..5b2cb9fe39b 100644 --- a/intern/cycles/blender/addon/engine.py +++ b/intern/cycles/blender/addon/engine.py @@ -242,7 +242,8 @@ def register_passes(engine, scene, srl): if crl.pass_debug_bvh_intersections: engine.register_pass(scene, srl, "Debug BVH Intersections", 1, "X", 'VALUE') if crl.pass_debug_ray_bounces: engine.register_pass(scene, srl, "Debug Ray Bounces", 1, "X", 'VALUE') - if crl.use_denoising and crl.denoising_store_passes: + cscene = scene.cycles + if crl.use_denoising and crl.denoising_store_passes and not cscene.use_progressive_refine: engine.register_pass(scene, srl, "Denoising Normal", 3, "XYZ", 'VECTOR') engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR') engine.register_pass(scene, srl, "Denoising Albedo", 3, "RGB", 'COLOR') diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index 28d60671bd8..309e44ccbb8 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -689,7 +689,11 @@ class CyclesRenderSettings(bpy.types.PropertyGroup): update=devices_update_callback ) - cls.debug_opencl_kernel_single_program = BoolProperty(name="Single Program", default=True, update=devices_update_callback); + cls.debug_opencl_kernel_single_program = BoolProperty( + name="Single Program", + default=True, + update=devices_update_callback, + ) cls.debug_use_opencl_debug = BoolProperty(name="Debug OpenCL", default=False) @@ -1203,6 +1207,7 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup): name="Use Denoising", description="Denoise the rendered image", default=False, + update=update_render_passes, ) cls.denoising_diffuse_direct = BoolProperty( name="Diffuse Direct", diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py index a8018e3824d..756d3b15a89 100644 --- a/intern/cycles/blender/addon/ui.py +++ b/intern/cycles/blender/addon/ui.py @@ -532,17 +532,17 @@ class CyclesRender_PT_layer_passes(CyclesButtonsPanel, Panel): col.prop(rl, "use_pass_environment") if context.scene.cycles.feature_set == 'EXPERIMENTAL': - col.separator() - sub = col.column() - sub.active = crl.use_denoising - sub.prop(crl, "denoising_store_passes", text="Denoising") + col.separator() + sub = col.column() + sub.active = crl.use_denoising + sub.prop(crl, "denoising_store_passes", text="Denoising") if _cycles.with_cycles_debug: - col = layout.column() - col.prop(crl, "pass_debug_bvh_traversed_nodes") - col.prop(crl, "pass_debug_bvh_traversed_instances") - col.prop(crl, "pass_debug_bvh_intersections") - col.prop(crl, "pass_debug_ray_bounces") + col = layout.column() + col.prop(crl, "pass_debug_bvh_traversed_nodes") + col.prop(crl, "pass_debug_bvh_traversed_instances") + col.prop(crl, "pass_debug_bvh_intersections") + col.prop(crl, "pass_debug_ray_bounces") class CyclesRender_PT_views(CyclesButtonsPanel, Panel): @@ -1688,7 +1688,7 @@ def draw_device(self, context): layout.prop(cscene, "feature_set") - split = layout.split(percentage=1/3) + split = layout.split(percentage=1 / 3) split.label("Device:") row = split.row() row.active = show_device_active(context) diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp index 46f32fe816d..d8c3750cf20 100644 --- a/intern/cycles/blender/blender_session.cpp +++ b/intern/cycles/blender/blender_session.cpp @@ -403,14 +403,7 @@ void BlenderSession::render() BL::RenderLayer b_rlay = *b_single_rlay; /* add passes */ - array<Pass> passes; - if(session_params.device.advanced_shading) { - passes = sync->sync_render_passes(b_rlay, *b_layer_iter); - } - else { - Pass::add(PASS_COMBINED, passes); - } - + array<Pass> passes = sync->sync_render_passes(b_rlay, *b_layer_iter, session_params); buffer_params.passes = passes; PointerRNA crl = RNA_pointer_get(&b_layer_iter->ptr, "cycles"); diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp index 1604b0039eb..ab986766211 100644 --- a/intern/cycles/blender/blender_sync.cpp +++ b/intern/cycles/blender/blender_sync.cpp @@ -537,11 +537,16 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass) } array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay, - BL::SceneRenderLayer& b_srlay) + BL::SceneRenderLayer& b_srlay, + const SessionParams &session_params) { array<Pass> passes; Pass::add(PASS_COMBINED, passes); + if(!session_params.device.advanced_shading) { + return passes; + } + /* loop over passes */ BL::RenderLayer::passes_iterator b_pass_iter; @@ -556,7 +561,9 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay, } PointerRNA crp = RNA_pointer_get(&b_srlay.ptr, "cycles"); - if(get_boolean(crp, "denoising_store_passes")) { + if(get_boolean(crp, "denoising_store_passes") && + get_boolean(crp, "use_denoising") && + !session_params.progressive_refine) { b_engine.add_pass("Denoising Normal", 3, "XYZ", b_srlay.name().c_str()); b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_srlay.name().c_str()); b_engine.add_pass("Denoising Albedo", 3, "RGB", b_srlay.name().c_str()); diff --git a/intern/cycles/blender/blender_sync.h b/intern/cycles/blender/blender_sync.h index a9f63346e00..fda8cb390c2 100644 --- a/intern/cycles/blender/blender_sync.h +++ b/intern/cycles/blender/blender_sync.h @@ -69,7 +69,8 @@ public: const char *layer = 0); void sync_render_layers(BL::SpaceView3D& b_v3d, const char *layer); array<Pass> sync_render_passes(BL::RenderLayer& b_rlay, - BL::SceneRenderLayer& b_srlay); + BL::SceneRenderLayer& b_srlay, + const SessionParams &session_params); void sync_integrator(); void sync_camera(BL::RenderSettings& b_render, BL::Object& b_override, diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp index 949c5f932a4..31671e76ec3 100644 --- a/intern/cycles/device/device.cpp +++ b/intern/cycles/device/device.cpp @@ -69,6 +69,8 @@ std::ostream& operator <<(std::ostream &os, << string_from_bool(requested_features.use_transparent) << std::endl; os << "Use Principled BSDF: " << string_from_bool(requested_features.use_principled) << std::endl; + os << "Use Denoising: " + << string_from_bool(requested_features.use_denoising) << std::endl; return os; } diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h index c22969d7dc6..68a555c1a93 100644 --- a/intern/cycles/device/device.h +++ b/intern/cycles/device/device.h @@ -127,6 +127,9 @@ public: /* Per-uber shader usage flags. */ bool use_principled; + /* Denoising features. */ + bool use_denoising; + DeviceRequestedFeatures() { /* TODO(sergey): Find more meaningful defaults. */ @@ -145,6 +148,7 @@ public: use_transparent = false; use_shadow_tricks = false; use_principled = false; + use_denoising = false; } bool modified(const DeviceRequestedFeatures& requested_features) @@ -163,7 +167,8 @@ public: use_patch_evaluation == requested_features.use_patch_evaluation && use_transparent == requested_features.use_transparent && use_shadow_tricks == requested_features.use_shadow_tricks && - use_principled == requested_features.use_principled); + use_principled == requested_features.use_principled && + use_denoising == requested_features.use_denoising); } /* Convert the requested features structure to a build options, @@ -213,6 +218,9 @@ public: if(!use_principled) { build_options += " -D__NO_PRINCIPLED__"; } + if(!use_denoising) { + build_options += " -D__NO_DENOISING__"; + } return build_options; } }; diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp index dddd19f179f..d2b3a89fa98 100644 --- a/intern/cycles/device/device_split_kernel.cpp +++ b/intern/cycles/device/device_split_kernel.cpp @@ -47,6 +47,7 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device) : device(device) kernel_direct_lighting = NULL; kernel_shadow_blocked_ao = NULL; kernel_shadow_blocked_dl = NULL; + kernel_enqueue_inactive = NULL; kernel_next_iteration_setup = NULL; kernel_indirect_subsurface = NULL; kernel_buffer_update = NULL; @@ -74,6 +75,7 @@ DeviceSplitKernel::~DeviceSplitKernel() delete kernel_direct_lighting; delete kernel_shadow_blocked_ao; delete kernel_shadow_blocked_dl; + delete kernel_enqueue_inactive; delete kernel_next_iteration_setup; delete kernel_indirect_subsurface; delete kernel_buffer_update; @@ -101,6 +103,7 @@ bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures& requested_fe LOAD_KERNEL(direct_lighting); LOAD_KERNEL(shadow_blocked_ao); LOAD_KERNEL(shadow_blocked_dl); + LOAD_KERNEL(enqueue_inactive); LOAD_KERNEL(next_iteration_setup); LOAD_KERNEL(indirect_subsurface); LOAD_KERNEL(buffer_update); @@ -256,6 +259,7 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task, ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size); ENQUEUE_SPLIT_KERNEL(shadow_blocked_ao, global_size, local_size); ENQUEUE_SPLIT_KERNEL(shadow_blocked_dl, global_size, local_size); + ENQUEUE_SPLIT_KERNEL(enqueue_inactive, global_size, local_size); ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size); ENQUEUE_SPLIT_KERNEL(indirect_subsurface, global_size, local_size); ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size); diff --git a/intern/cycles/device/device_split_kernel.h b/intern/cycles/device/device_split_kernel.h index 68c2ba974a5..2bac1998cb7 100644 --- a/intern/cycles/device/device_split_kernel.h +++ b/intern/cycles/device/device_split_kernel.h @@ -69,6 +69,7 @@ private: SplitKernelFunction *kernel_direct_lighting; SplitKernelFunction *kernel_shadow_blocked_ao; SplitKernelFunction *kernel_shadow_blocked_dl; + SplitKernelFunction *kernel_enqueue_inactive; SplitKernelFunction *kernel_next_iteration_setup; SplitKernelFunction *kernel_indirect_subsurface; SplitKernelFunction *kernel_buffer_update; diff --git a/intern/cycles/device/opencl/opencl.h b/intern/cycles/device/opencl/opencl.h index 52851061d7b..399fae9b42e 100644 --- a/intern/cycles/device/opencl/opencl.h +++ b/intern/cycles/device/opencl/opencl.h @@ -130,6 +130,11 @@ public: cl_int* error = NULL); static cl_device_type get_device_type(cl_device_id device_id); + static bool get_driver_version(cl_device_id device_id, + int *major, + int *minor, + cl_int* error = NULL); + static int mem_address_alignment(cl_device_id device_id); /* Get somewhat more readable device name. diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index 76dcbd6fc9a..08b632ee9d3 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -176,17 +176,62 @@ protected: friend class OpenCLSplitKernelFunction; }; +struct CachedSplitMemory { + int id; + device_memory *split_data; + device_memory *ray_state; + device_ptr *rng_state; + device_memory *queue_index; + device_memory *use_queues_flag; + device_memory *work_pools; + device_ptr *buffer; +}; + class OpenCLSplitKernelFunction : public SplitKernelFunction { public: OpenCLDeviceSplitKernel* device; OpenCLDeviceBase::OpenCLProgram program; + CachedSplitMemory& cached_memory; + int cached_id; + + OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device, CachedSplitMemory& cached_memory) : + device(device), cached_memory(cached_memory), cached_id(cached_memory.id-1) + { + } - OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device) : device(device) {} - ~OpenCLSplitKernelFunction() { program.release(); } + ~OpenCLSplitKernelFunction() + { + program.release(); + } virtual bool enqueue(const KernelDimensions& dim, device_memory& kg, device_memory& data) { - device->kernel_set_args(program(), 0, kg, data); + if(cached_id != cached_memory.id) { + cl_uint start_arg_index = + device->kernel_set_args(program(), + 0, + kg, + data, + *cached_memory.split_data, + *cached_memory.ray_state, + *cached_memory.rng_state); + +/* TODO(sergey): Avoid map lookup here. */ +#define KERNEL_TEX(type, ttype, name) \ + device->set_kernel_arg_mem(program(), &start_arg_index, #name); +#include "kernel/kernel_textures.h" +#undef KERNEL_TEX + + start_arg_index += + device->kernel_set_args(program(), + start_arg_index, + *cached_memory.queue_index, + *cached_memory.use_queues_flag, + *cached_memory.work_pools, + *cached_memory.buffer); + + cached_id = cached_memory.id; + } device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue, program(), @@ -213,6 +258,7 @@ public: class OpenCLSplitKernel : public DeviceSplitKernel { OpenCLDeviceSplitKernel *device; + CachedSplitMemory cached_memory; public: explicit OpenCLSplitKernel(OpenCLDeviceSplitKernel *device) : DeviceSplitKernel(device), device(device) { } @@ -220,7 +266,7 @@ public: virtual SplitKernelFunction* get_split_kernel_function(string kernel_name, const DeviceRequestedFeatures& requested_features) { - OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device); + OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device, cached_memory); bool single_program = OpenCLInfo::use_single_program(); kernel->program = @@ -349,6 +395,15 @@ public: return false; } + cached_memory.split_data = &split_data; + cached_memory.ray_state = &ray_state; + cached_memory.rng_state = &rtile.rng_state; + cached_memory.queue_index = &queue_index; + cached_memory.use_queues_flag = &use_queues_flag; + cached_memory.work_pools = &work_pool_wgs; + cached_memory.buffer = &rtile.buffer; + cached_memory.id++; + return true; } diff --git a/intern/cycles/device/opencl/opencl_util.cpp b/intern/cycles/device/opencl/opencl_util.cpp index d27aa05c312..8ba2a8e26da 100644 --- a/intern/cycles/device/opencl/opencl_util.cpp +++ b/intern/cycles/device/opencl/opencl_util.cpp @@ -608,6 +608,14 @@ bool OpenCLInfo::device_supported(const string& platform_name, if(!get_device_name(device_id, &device_name)) { return false; } + + int driver_major = 0; + int driver_minor = 0; + if(!get_driver_version(device_id, &driver_major, &driver_minor)) { + return false; + } + VLOG(3) << "OpenCL driver version " << driver_major << "." << driver_minor; + /* It is possible tyo have Iris GPU on AMD/Apple OpenCL framework * (aka, it will not be on Intel framework). This isn't supported * and needs an explicit blacklist. @@ -618,6 +626,21 @@ bool OpenCLInfo::device_supported(const string& platform_name, if(platform_name == "AMD Accelerated Parallel Processing" && device_type == CL_DEVICE_TYPE_GPU) { + if(driver_major < 2236) { + VLOG(1) << "AMD driver version " << driver_major << "." << driver_minor << " not supported."; + return false; + } + const char *blacklist[] = { + /* GCN 1 */ + "Tahiti", "Pitcairn", "Capeverde", "Oland", + NULL + }; + for (int i = 0; blacklist[i] != NULL; i++) { + if(device_name == blacklist[i]) { + VLOG(1) << "AMD device " << device_name << " not supported"; + return false; + } + } return true; } if(platform_name == "Apple" && device_type == CL_DEVICE_TYPE_GPU) { @@ -1073,6 +1096,34 @@ string OpenCLInfo::get_readable_device_name(cl_device_id device_id) return get_device_name(device_id); } +bool OpenCLInfo::get_driver_version(cl_device_id device_id, + int *major, + int *minor, + cl_int* error) +{ + char buffer[1024]; + cl_int err; + if((err = clGetDeviceInfo(device_id, + CL_DRIVER_VERSION, + sizeof(buffer), + &buffer, + NULL)) != CL_SUCCESS) + { + if(error != NULL) { + *error = err; + } + return false; + } + if(error != NULL) { + *error = CL_SUCCESS; + } + if(sscanf(buffer, "%d.%d", major, minor) < 2) { + VLOG(1) << string_printf("OpenCL: failed to parse driver version string (%s).", buffer); + return false; + } + return true; +} + int OpenCLInfo::mem_address_alignment(cl_device_id device_id) { int base_align_bits; diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index bef869f34b4..23e9bd311c4 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -45,6 +45,7 @@ set(SRC kernels/opencl/kernel_direct_lighting.cl kernels/opencl/kernel_shadow_blocked_ao.cl kernels/opencl/kernel_shadow_blocked_dl.cl + kernels/opencl/kernel_enqueue_inactive.cl kernels/opencl/kernel_next_iteration_setup.cl kernels/opencl/kernel_indirect_subsurface.cl kernels/opencl/kernel_buffer_update.cl @@ -121,6 +122,10 @@ set(SRC_KERNELS_CUDA_HEADERS kernels/cuda/kernel_config.h ) +set(SRC_KERNELS_OPENCL_HEADERS + kernels/opencl/kernel_split_function.h +) + set(SRC_CLOSURE_HEADERS closure/alloc.h closure/bsdf.h @@ -278,6 +283,7 @@ set(SRC_SPLIT_HEADERS split/kernel_data_init.h split/kernel_direct_lighting.h split/kernel_do_volume.h + split/kernel_enqueue_inactive.h split/kernel_holdout_emission_blurring_pathtermination_ao.h split/kernel_indirect_background.h split/kernel_indirect_subsurface.h @@ -450,6 +456,7 @@ add_library(cycles_kernel ${SRC_HEADERS} ${SRC_KERNELS_CPU_HEADERS} ${SRC_KERNELS_CUDA_HEADERS} + ${SRC_KERNELS_OPENCL_HEADERS} ${SRC_BVH_HEADERS} ${SRC_CLOSURE_HEADERS} ${SRC_FILTER_HEADERS} @@ -490,9 +497,11 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_subsurface_sc delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_direct_lighting.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_ao.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_dl.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) +delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_enqueue_inactive.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_next_iteration_setup.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_indirect_subsurface.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_buffer_update.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) +delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_split_function.h" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/filter.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel_split.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda) diff --git a/intern/cycles/kernel/filter/filter_nlm_cpu.h b/intern/cycles/kernel/filter/filter_nlm_cpu.h index 88afc00ccb3..3e752bce68f 100644 --- a/intern/cycles/kernel/filter/filter_nlm_cpu.h +++ b/intern/cycles/kernel/filter/filter_nlm_cpu.h @@ -101,7 +101,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict d for(int x = rect.x; x < rect.z; x++) { const int low = max(rect.x, x-f); const int high = min(rect.z, x+f+1); - out_image[y*w+x] = expf(-max(out_image[y*w+x] * (1.0f/(high - low)), 0.0f)); + out_image[y*w+x] = fast_expf(-max(out_image[y*w+x] * (1.0f/(high - low)), 0.0f)); } } } diff --git a/intern/cycles/kernel/filter/filter_nlm_gpu.h b/intern/cycles/kernel/filter/filter_nlm_gpu.h index 62bd5be1de5..2c5ac807051 100644 --- a/intern/cycles/kernel/filter/filter_nlm_gpu.h +++ b/intern/cycles/kernel/filter/filter_nlm_gpu.h @@ -66,7 +66,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(int x, int y, sum += difference_image[y*w+x1]; } sum *= 1.0f/(high-low); - out_image[y*w+x] = expf(-max(sum, 0.0f)); + out_image[y*w+x] = fast_expf(-max(sum, 0.0f)); } ccl_device_inline void kernel_filter_nlm_update_output(int x, int y, diff --git a/intern/cycles/kernel/filter/filter_reconstruction.h b/intern/cycles/kernel/filter/filter_reconstruction.h index 90a2816ddf7..25a3025056c 100644 --- a/intern/cycles/kernel/filter/filter_reconstruction.h +++ b/intern/cycles/kernel/filter/filter_reconstruction.h @@ -29,20 +29,24 @@ ccl_device_inline void kernel_filter_construct_gramian(int x, int y, ccl_global float3 *XtWY, int localIdx) { + if(weight < 1e-3f) { + return; + } + int p_offset = y *w + x; int q_offset = (y+dy)*w + (x+dx); -#ifdef __KERNEL_CPU__ - const int stride = 1; - (void)storage_stride; - (void)localIdx; - float design_row[DENOISE_FEATURES+1]; -#elif defined(__KERNEL_CUDA__) +#ifdef __KERNEL_GPU__ const int stride = storage_stride; +#else + const int stride = 1; + (void) storage_stride; +#endif + +#ifdef __KERNEL_CUDA__ ccl_local float shared_design_row[(DENOISE_FEATURES+1)*CCL_MAX_LOCAL_SIZE]; ccl_local_param float *design_row = shared_design_row + localIdx*(DENOISE_FEATURES+1); #else - const int stride = storage_stride; float design_row[DENOISE_FEATURES+1]; #endif @@ -70,13 +74,19 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h, int4 buffer_params, int sample) { -#ifdef __KERNEL_CPU__ - const int stride = 1; - (void)storage_stride; -#else +#ifdef __KERNEL_GPU__ const int stride = storage_stride; +#else + const int stride = 1; + (void) storage_stride; #endif + if(XtWX[0] < 1e-3f) { + /* There is not enough information to determine a denoised result. + * As a fallback, keep the original value of the pixel. */ + return; + } + /* The weighted average of pixel colors (essentially, the NLM-filtered image). * In case the solution of the linear model fails due to numerical issues, * fall back to this value. */ @@ -89,6 +99,9 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h, final_color = mean_color; } + /* Clamp pixel value to positive values. */ + final_color = max(final_color, make_float3(0.0f, 0.0f, 0.0f)); + ccl_global float *combined_buffer = buffer + (y*buffer_params.y + x + buffer_params.x)*buffer_params.z; final_color *= sample; if(buffer_params.w) { @@ -101,6 +114,4 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h, combined_buffer[2] = final_color.z; } -#undef STORAGE_TYPE - CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h index 06728415c15..175bd6b9737 100644 --- a/intern/cycles/kernel/kernel_accumulate.h +++ b/intern/cycles/kernel/kernel_accumulate.h @@ -621,25 +621,43 @@ ccl_device_inline void path_radiance_accum_sample(PathRadiance *L, PathRadiance { float fac = 1.0f/num_samples; +#ifdef __SPLIT_KERNEL__ +# define safe_float3_add(f, v) \ + do { \ + ccl_global float *p = (ccl_global float*)(&(f)); \ + atomic_add_and_fetch_float(p+0, (v).x); \ + atomic_add_and_fetch_float(p+1, (v).y); \ + atomic_add_and_fetch_float(p+2, (v).z); \ + } while(0) +#else +# define safe_float3_add(f, v) (f) += (v) +#endif /* __SPLIT_KERNEL__ */ + #ifdef __PASSES__ - L->direct_diffuse += L_sample->direct_diffuse*fac; - L->direct_glossy += L_sample->direct_glossy*fac; - L->direct_transmission += L_sample->direct_transmission*fac; - L->direct_subsurface += L_sample->direct_subsurface*fac; - L->direct_scatter += L_sample->direct_scatter*fac; - - L->indirect_diffuse += L_sample->indirect_diffuse*fac; - L->indirect_glossy += L_sample->indirect_glossy*fac; - L->indirect_transmission += L_sample->indirect_transmission*fac; - L->indirect_subsurface += L_sample->indirect_subsurface*fac; - L->indirect_scatter += L_sample->indirect_scatter*fac; - - L->background += L_sample->background*fac; - L->ao += L_sample->ao*fac; - L->shadow += L_sample->shadow*fac; + safe_float3_add(L->direct_diffuse, L_sample->direct_diffuse*fac); + safe_float3_add(L->direct_glossy, L_sample->direct_glossy*fac); + safe_float3_add(L->direct_transmission, L_sample->direct_transmission*fac); + safe_float3_add(L->direct_subsurface, L_sample->direct_subsurface*fac); + safe_float3_add(L->direct_scatter, L_sample->direct_scatter*fac); + + safe_float3_add(L->indirect_diffuse, L_sample->indirect_diffuse*fac); + safe_float3_add(L->indirect_glossy, L_sample->indirect_glossy*fac); + safe_float3_add(L->indirect_transmission, L_sample->indirect_transmission*fac); + safe_float3_add(L->indirect_subsurface, L_sample->indirect_subsurface*fac); + safe_float3_add(L->indirect_scatter, L_sample->indirect_scatter*fac); + + safe_float3_add(L->background, L_sample->background*fac); + safe_float3_add(L->ao, L_sample->ao*fac); + safe_float3_add(L->shadow, L_sample->shadow*fac); +# ifdef __SPLIT_KERNEL__ + atomic_add_and_fetch_float(&L->mist, L_sample->mist*fac); +# else L->mist += L_sample->mist*fac; -#endif - L->emission += L_sample->emission * fac; +# endif /* __SPLIT_KERNEL__ */ +#endif /* __PASSES__ */ + safe_float3_add(L->emission, L_sample->emission*fac); + +#undef safe_float3_add } #ifdef __SHADOW_TRICKS__ diff --git a/intern/cycles/kernel/kernel_path_state.h b/intern/cycles/kernel/kernel_path_state.h index 0fa77d9e8bd..5d92fd12201 100644 --- a/intern/cycles/kernel/kernel_path_state.h +++ b/intern/cycles/kernel/kernel_path_state.h @@ -139,9 +139,11 @@ ccl_device_inline void path_state_next(KernelGlobals *kg, ccl_addr_space PathSta /* random number generator next bounce */ state->rng_offset += PRNG_BOUNCE_NUM; +#ifdef __DENOISING_FEATURES__ if((state->denoising_feature_weight == 0.0f) && !(state->flag & PATH_RAY_SHADOW_CATCHER)) { state->flag &= ~PATH_RAY_STORE_SHADOW_INFO; } +#endif } ccl_device_inline uint path_state_ray_visibility(KernelGlobals *kg, PathState *state) diff --git a/intern/cycles/kernel/kernel_queues.h b/intern/cycles/kernel/kernel_queues.h index 96bc636d5ac..e32d4bbbc1b 100644 --- a/intern/cycles/kernel/kernel_queues.h +++ b/intern/cycles/kernel/kernel_queues.h @@ -128,6 +128,21 @@ ccl_device unsigned int get_global_queue_index( return my_gqidx; } +ccl_device int dequeue_ray_index( + int queue_number, + ccl_global int *queues, + int queue_size, + ccl_global int *queue_index) +{ + int index = atomic_fetch_and_dec_uint32((ccl_global uint*)&queue_index[queue_number])-1; + + if(index < 0) { + return QUEUE_EMPTY_SLOT; + } + + return queues[index + queue_number * queue_size]; +} + CCL_NAMESPACE_END #endif // __KERNEL_QUEUE_H__ diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index dbeaffdfb24..31e47e837fd 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -236,6 +236,9 @@ CCL_NAMESPACE_BEGIN #ifdef __NO_PRINCIPLED__ # undef __PRINCIPLED__ #endif +#ifdef __NO_DENOISING__ +# undef __DENOISING_FEATURES__ +#endif /* Random Numbers */ @@ -1387,6 +1390,8 @@ enum QueueNumber { #ifdef __BRANCHED_PATH__ /* All rays moving to next iteration of the indirect loop for light */ QUEUE_LIGHT_INDIRECT_ITER, + /* Queue of all inactive rays. These are candidates for sharing work of indirect loops */ + QUEUE_INACTIVE_RAYS, # ifdef __VOLUME__ /* All rays moving to next iteration of the indirect loop for volumes */ QUEUE_VOLUME_INDIRECT_ITER, @@ -1429,6 +1434,9 @@ enum RayState { RAY_BRANCHED_VOLUME_INDIRECT = (1 << 5), RAY_BRANCHED_SUBSURFACE_INDIRECT = (1 << 6), RAY_BRANCHED_INDIRECT = (RAY_BRANCHED_LIGHT_INDIRECT | RAY_BRANCHED_VOLUME_INDIRECT | RAY_BRANCHED_SUBSURFACE_INDIRECT), + + /* Ray is evaluating an iteration of an indirect loop for another thread */ + RAY_BRANCHED_INDIRECT_SHARED = (1 << 7), }; #define ASSIGN_RAY_STATE(ray_state, ray_index, state) (ray_state[ray_index] = ((ray_state[ray_index] & RAY_FLAG_MASK) | state)) diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h index 9895080d328..c8938534fe8 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h @@ -85,6 +85,7 @@ DECLARE_SPLIT_KERNEL_FUNCTION(subsurface_scatter) DECLARE_SPLIT_KERNEL_FUNCTION(direct_lighting) DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao) DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl) +DECLARE_SPLIT_KERNEL_FUNCTION(enqueue_inactive) DECLARE_SPLIT_KERNEL_FUNCTION(next_iteration_setup) DECLARE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DECLARE_SPLIT_KERNEL_FUNCTION(buffer_update) diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h index 9b85a864153..d4315ee5ec4 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h @@ -53,6 +53,7 @@ # include "kernel/split/kernel_direct_lighting.h" # include "kernel/split/kernel_shadow_blocked_ao.h" # include "kernel/split/kernel_shadow_blocked_dl.h" +# include "kernel/split/kernel_enqueue_inactive.h" # include "kernel/split/kernel_next_iteration_setup.h" # include "kernel/split/kernel_indirect_subsurface.h" # include "kernel/split/kernel_buffer_update.h" @@ -230,6 +231,7 @@ DEFINE_SPLIT_KERNEL_FUNCTION(subsurface_scatter) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(direct_lighting, uint) DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao) DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl) +DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint) DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint) diff --git a/intern/cycles/kernel/kernels/cuda/kernel_split.cu b/intern/cycles/kernel/kernels/cuda/kernel_split.cu index 8b7f1a8d405..628891b1458 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel_split.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel_split.cu @@ -39,6 +39,7 @@ #include "kernel/split/kernel_direct_lighting.h" #include "kernel/split/kernel_shadow_blocked_ao.h" #include "kernel/split/kernel_shadow_blocked_dl.h" +#include "kernel/split/kernel_enqueue_inactive.h" #include "kernel/split/kernel_next_iteration_setup.h" #include "kernel/split/kernel_indirect_subsurface.h" #include "kernel/split/kernel_buffer_update.h" @@ -118,6 +119,7 @@ DEFINE_SPLIT_KERNEL_FUNCTION(subsurface_scatter) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(direct_lighting, uint) DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao) DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl) +DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint) DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint) diff --git a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl index db65c91baf7..dcea2630aef 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl @@ -18,10 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_buffer_update.h" -__kernel void kernel_ocl_path_trace_buffer_update( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local unsigned int local_queue_atomics; - kernel_buffer_update((KernelGlobals*)kg, &local_queue_atomics); -} +#define KERNEL_NAME buffer_update +#define LOCALS_TYPE unsigned int +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl index eb34f750881..ed64ae01aae 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl @@ -18,10 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_direct_lighting.h" -__kernel void kernel_ocl_path_trace_direct_lighting( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local unsigned int local_queue_atomics; - kernel_direct_lighting((KernelGlobals*)kg, &local_queue_atomics); -} +#define KERNEL_NAME direct_lighting +#define LOCALS_TYPE unsigned int +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl index 83ef5f5f3f2..8afaa686e28 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_do_volume.h" -__kernel void kernel_ocl_path_trace_do_volume( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_do_volume((KernelGlobals*)kg); -} +#define KERNEL_NAME do_volume +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl b/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl new file mode 100644 index 00000000000..e68d4104a91 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl @@ -0,0 +1,26 @@ +/* + * Copyright 2011-2017 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/split/kernel_split_common.h" +#include "kernel/split/kernel_enqueue_inactive.h" + +#define KERNEL_NAME enqueue_inactive +#define LOCALS_TYPE unsigned int +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl index d071b39aa6f..9e1e57beba6 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl @@ -18,12 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h" -__kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local BackgroundAOLocals locals; - kernel_holdout_emission_blurring_pathtermination_ao( - (KernelGlobals*)kg, - &locals); -} +#define KERNEL_NAME holdout_emission_blurring_pathtermination_ao +#define LOCALS_TYPE BackgroundAOLocals +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl index 8c213ff5cb2..192d01444ba 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_indirect_background.h" -__kernel void kernel_ocl_path_trace_indirect_background( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_indirect_background((KernelGlobals*)kg); -} +#define KERNEL_NAME indirect_background +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl index 998ebc4c0c3..84938b889e5 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_indirect_subsurface.h" -__kernel void kernel_ocl_path_trace_indirect_subsurface( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_indirect_subsurface((KernelGlobals*)kg); -} +#define KERNEL_NAME indirect_subsurface +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl index 822d2287715..c314dc96c33 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_lamp_emission.h" -__kernel void kernel_ocl_path_trace_lamp_emission( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_lamp_emission((KernelGlobals*)kg); -} +#define KERNEL_NAME lamp_emission +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl index 6d207253a40..8b1332bf013 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl @@ -18,10 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_next_iteration_setup.h" -__kernel void kernel_ocl_path_trace_next_iteration_setup( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local unsigned int local_queue_atomics; - kernel_next_iteration_setup((KernelGlobals*)kg, &local_queue_atomics); -} +#define KERNEL_NAME next_iteration_setup +#define LOCALS_TYPE unsigned int +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl index bd9aa9538c8..fa210e747c0 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_path_init.h" -__kernel void kernel_ocl_path_trace_path_init( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_path_init((KernelGlobals*)kg); -} +#define KERNEL_NAME path_init +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl index 9be154e3d75..68ee6f1d536 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl @@ -18,10 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_queue_enqueue.h" -__kernel void kernel_ocl_path_trace_queue_enqueue( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local QueueEnqueueLocals locals; - kernel_queue_enqueue((KernelGlobals*)kg, &locals); -} +#define KERNEL_NAME queue_enqueue +#define LOCALS_TYPE QueueEnqueueLocals +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl index eb4fb4d153a..10d09377ba9 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_scene_intersect.h" -__kernel void kernel_ocl_path_trace_scene_intersect( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_scene_intersect((KernelGlobals*)kg); -} +#define KERNEL_NAME scene_intersect +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl index 5bfb31b193a..40eaa561863 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_shader_eval.h" -__kernel void kernel_ocl_path_trace_shader_eval( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_shader_eval((KernelGlobals*)kg); -} +#define KERNEL_NAME shader_eval +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl index 38bfd04ad4c..8c36100f762 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl @@ -18,10 +18,9 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_shader_setup.h" -__kernel void kernel_ocl_path_trace_shader_setup( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local unsigned int local_queue_atomics; - kernel_shader_setup((KernelGlobals*)kg, &local_queue_atomics); -} +#define KERNEL_NAME shader_setup +#define LOCALS_TYPE unsigned int +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl index 6f722915d45..bcacaa4a054 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl @@ -19,10 +19,9 @@ #include "kernel/split/kernel_shader_sort.h" __attribute__((reqd_work_group_size(64, 1, 1))) -__kernel void kernel_ocl_path_trace_shader_sort( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - ccl_local ShaderSortLocals locals; - kernel_shader_sort((KernelGlobals*)kg, &locals); -} +#define KERNEL_NAME shader_sort +#define LOCALS_TYPE ShaderSortLocals +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME +#undef LOCALS_TYPE + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl index 6a8ef81b32a..8de250a375c 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_shadow_blocked_ao.h" -__kernel void kernel_ocl_path_trace_shadow_blocked_ao( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_shadow_blocked_ao((KernelGlobals*)kg); -} +#define KERNEL_NAME shadow_blocked_ao +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl index b255cc5ef8b..29da77022ed 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_shadow_blocked_dl.h" -__kernel void kernel_ocl_path_trace_shadow_blocked_dl( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_shadow_blocked_dl((KernelGlobals*)kg); -} +#define KERNEL_NAME shadow_blocked_dl +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split.cl b/intern/cycles/kernel/kernels/opencl/kernel_split.cl index 8de82db7afe..651addb02f4 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_split.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_split.cl @@ -31,6 +31,7 @@ #include "kernel/kernels/opencl/kernel_direct_lighting.cl" #include "kernel/kernels/opencl/kernel_shadow_blocked_ao.cl" #include "kernel/kernels/opencl/kernel_shadow_blocked_dl.cl" +#include "kernel/kernels/opencl/kernel_enqueue_inactive.cl" #include "kernel/kernels/opencl/kernel_next_iteration_setup.cl" #include "kernel/kernels/opencl/kernel_indirect_subsurface.cl" #include "kernel/kernels/opencl/kernel_buffer_update.cl" diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h new file mode 100644 index 00000000000..f1e914a70d4 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h @@ -0,0 +1,72 @@ +/* + * Copyright 2011-2017 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#define KERNEL_NAME_JOIN(a, b) a ## _ ## b +#define KERNEL_NAME_EVAL(a, b) KERNEL_NAME_JOIN(a, b) + +__kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)( + ccl_global char *kg_global, + ccl_constant KernelData *data, + + ccl_global void *split_data_buffer, + ccl_global char *ray_state, + ccl_global uint *rng_state, + +#define KERNEL_TEX(type, ttype, name) \ + ccl_global type *name, +#include "kernel/kernel_textures.h" + + ccl_global int *queue_index, + ccl_global char *use_queues_flag, + ccl_global unsigned int *work_pools, + ccl_global float *buffer + ) +{ +#ifdef LOCALS_TYPE + ccl_local LOCALS_TYPE locals; +#endif + + KernelGlobals *kg = (KernelGlobals*)kg_global; + + if(ccl_local_id(0) + ccl_local_id(1) == 0) { + kg->data = data; + + kernel_split_params.rng_state = rng_state; + kernel_split_params.queue_index = queue_index; + kernel_split_params.use_queues_flag = use_queues_flag; + kernel_split_params.work_pools = work_pools; + kernel_split_params.buffer = buffer; + + split_data_init(kg, &kernel_split_state, ccl_global_size(0)*ccl_global_size(1), split_data_buffer, ray_state); + +#define KERNEL_TEX(type, ttype, name) \ + kg->name = name; +#include "kernel/kernel_textures.h" + } + + ccl_barrier(CCL_LOCAL_MEM_FENCE); + + KERNEL_NAME_EVAL(kernel, KERNEL_NAME)( + kg +#ifdef LOCALS_TYPE + , &locals +#endif + ); +} + +#undef KERNEL_NAME_JOIN +#undef KERNEL_NAME_EVAL + diff --git a/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl b/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl index 99b74a1802b..2b3be38df84 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_subsurface_scatter.h" -__kernel void kernel_ocl_path_trace_subsurface_scatter( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_subsurface_scatter((KernelGlobals*)kg); -} +#define KERNEL_NAME subsurface_scatter +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + diff --git a/intern/cycles/kernel/split/kernel_branched.h b/intern/cycles/kernel/split/kernel_branched.h index dc74a2ada53..e2762a85fc8 100644 --- a/intern/cycles/kernel/split/kernel_branched.h +++ b/intern/cycles/kernel/split/kernel_branched.h @@ -63,12 +63,49 @@ ccl_device_inline void kernel_split_branched_path_indirect_loop_end(KernelGlobal REMOVE_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_INDIRECT); } +ccl_device_inline bool kernel_split_branched_indirect_start_shared(KernelGlobals *kg, int ray_index) +{ + ccl_global char *ray_state = kernel_split_state.ray_state; + + int inactive_ray = dequeue_ray_index(QUEUE_INACTIVE_RAYS, + kernel_split_state.queue_data, kernel_split_params.queue_size, kernel_split_params.queue_index); + + if(!IS_STATE(ray_state, inactive_ray, RAY_INACTIVE)) { + return false; + } + +#define SPLIT_DATA_ENTRY(type, name, num) \ + kernel_split_state.name[inactive_ray] = kernel_split_state.name[ray_index]; + SPLIT_DATA_ENTRIES_BRANCHED_SHARED +#undef SPLIT_DATA_ENTRY + + kernel_split_state.branched_state[inactive_ray].shared_sample_count = 0; + kernel_split_state.branched_state[inactive_ray].original_ray = ray_index; + kernel_split_state.branched_state[inactive_ray].waiting_on_shared_samples = false; + + PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; + PathRadiance *inactive_L = &kernel_split_state.path_radiance[inactive_ray]; + + path_radiance_init(inactive_L, kernel_data.film.use_light_pass); + inactive_L->direct_throughput = L->direct_throughput; + path_radiance_copy_indirect(inactive_L, L); + + ray_state[inactive_ray] = RAY_REGENERATED; + ADD_RAY_FLAG(ray_state, inactive_ray, RAY_BRANCHED_INDIRECT_SHARED); + ADD_RAY_FLAG(ray_state, inactive_ray, IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)); + + atomic_fetch_and_inc_uint32((ccl_global uint*)&kernel_split_state.branched_state[ray_index].shared_sample_count); + + return true; +} + /* bounce off surface and integrate indirect light */ ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter(KernelGlobals *kg, int ray_index, float num_samples_adjust, ShaderData *saved_sd, - bool reset_path_state) + bool reset_path_state, + bool wait_for_shared) { SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; @@ -155,12 +192,25 @@ ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter( /* start the indirect path */ *tp *= num_samples_inv; + if(kernel_split_branched_indirect_start_shared(kg, ray_index)) { + continue; + } + return true; } branched_state->next_sample = 0; } + branched_state->next_closure = sd->num_closure; + + if(wait_for_shared) { + branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); + if(branched_state->waiting_on_shared_samples) { + return true; + } + } + return false; } diff --git a/intern/cycles/kernel/split/kernel_do_volume.h b/intern/cycles/kernel/split/kernel_do_volume.h index 694b777f429..9f8dd2392d9 100644 --- a/intern/cycles/kernel/split/kernel_do_volume.h +++ b/intern/cycles/kernel/split/kernel_do_volume.h @@ -75,11 +75,30 @@ ccl_device_noinline bool kernel_split_branched_path_volume_indirect_light_iter(K branched_state->next_sample = j+1; branched_state->num_samples = num_samples; + /* Attempting to share too many samples is slow for volumes as it causes us to + * loop here more and have many calls to kernel_volume_integrate which evaluates + * shaders. The many expensive shader evaluations cause the work load to become + * unbalanced and many threads to become idle in this kernel. Limiting the + * number of shared samples here helps quite a lot. + */ + if(branched_state->shared_sample_count < 2) { + if(kernel_split_branched_indirect_start_shared(kg, ray_index)) { + continue; + } + } + return true; } # endif } + branched_state->next_sample = num_samples; + + branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); + if(branched_state->waiting_on_shared_samples) { + return true; + } + kernel_split_branched_path_indirect_loop_end(kg, ray_index); /* todo: avoid this calculation using decoupled ray marching */ diff --git a/intern/cycles/kernel/split/kernel_enqueue_inactive.h b/intern/cycles/kernel/split/kernel_enqueue_inactive.h new file mode 100644 index 00000000000..496355bbc3a --- /dev/null +++ b/intern/cycles/kernel/split/kernel_enqueue_inactive.h @@ -0,0 +1,46 @@ +/* + * Copyright 2011-2017 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +CCL_NAMESPACE_BEGIN + +ccl_device void kernel_enqueue_inactive(KernelGlobals *kg, + ccl_local_param unsigned int *local_queue_atomics) +{ +#ifdef __BRANCHED_PATH__ + /* Enqeueue RAY_INACTIVE rays into QUEUE_INACTIVE_RAYS queue. */ + if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { + *local_queue_atomics = 0; + } + ccl_barrier(CCL_LOCAL_MEM_FENCE); + + int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); + + char enqueue_flag = 0; + if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_INACTIVE)) { + enqueue_flag = 1; + } + + enqueue_ray_index_local(ray_index, + QUEUE_INACTIVE_RAYS, + enqueue_flag, + kernel_split_params.queue_size, + local_queue_atomics, + kernel_split_state.queue_data, + kernel_split_params.queue_index); +#endif /* __BRANCHED_PATH__ */ +} + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h index 71017fed19e..7758e35fd32 100644 --- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h +++ b/intern/cycles/kernel/split/kernel_next_iteration_setup.h @@ -147,6 +147,7 @@ ccl_device void kernel_next_iteration_setup(KernelGlobals *kg, ray_index, 1.0f, &kernel_split_state.branched_state[ray_index].sd, + true, true)) { ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); @@ -193,6 +194,7 @@ ccl_device void kernel_next_iteration_setup(KernelGlobals *kg, ray_index, 1.0f, &kernel_split_state.branched_state[ray_index].sd, + true, true)) { ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h index 5dc94caec85..45984ca509b 100644 --- a/intern/cycles/kernel/split/kernel_scene_intersect.h +++ b/intern/cycles/kernel/split/kernel_scene_intersect.h @@ -43,11 +43,21 @@ ccl_device void kernel_scene_intersect(KernelGlobals *kg) } /* All regenerated rays become active here */ - if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) - ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE); + if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) { +#ifdef __BRANCHED_PATH__ + if(kernel_split_state.branched_state[ray_index].waiting_on_shared_samples) { + kernel_split_path_end(kg, ray_index); + } + else +#endif /* __BRANCHED_PATH__ */ + { + ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE); + } + } - if(!IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) + if(!IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) { return; + } #ifdef __KERNEL_DEBUG__ DebugData *debug_data = &kernel_split_state.debug_data[ray_index]; diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h b/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h index 386fbbc4d09..78e61709b01 100644 --- a/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h +++ b/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h @@ -29,6 +29,14 @@ ccl_device void kernel_shadow_blocked_dl(KernelGlobals *kg) kernel_split_state.queue_data, kernel_split_params.queue_size, 1); } +#ifdef __BRANCHED_PATH__ + /* TODO(mai): move this somewhere else? */ + if(thread_index == 0) { + /* Clear QUEUE_INACTIVE_RAYS before next kernel. */ + kernel_split_params.queue_index[QUEUE_INACTIVE_RAYS] = 0; + } +#endif /* __BRANCHED_PATH__ */ + if(ray_index == QUEUE_EMPTY_SLOT) return; diff --git a/intern/cycles/kernel/split/kernel_split_common.h b/intern/cycles/kernel/split/kernel_split_common.h index 57f070d51e0..08f0124b529 100644 --- a/intern/cycles/kernel/split/kernel_split_common.h +++ b/intern/cycles/kernel/split/kernel_split_common.h @@ -56,7 +56,20 @@ ccl_device_inline void kernel_split_path_end(KernelGlobals *kg, int ray_index) ccl_global char *ray_state = kernel_split_state.ray_state; #ifdef __BRANCHED_PATH__ - if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT)) { + if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT_SHARED)) { + int orig_ray = kernel_split_state.branched_state[ray_index].original_ray; + + PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; + PathRadiance *orig_ray_L = &kernel_split_state.path_radiance[orig_ray]; + + path_radiance_sum_indirect(L); + path_radiance_accum_sample(orig_ray_L, L, 1); + + atomic_fetch_and_dec_uint32((ccl_global uint*)&kernel_split_state.branched_state[orig_ray].shared_sample_count); + + ASSIGN_RAY_STATE(ray_state, ray_index, RAY_INACTIVE); + } + else if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT)) { ASSIGN_RAY_STATE(ray_state, ray_index, RAY_LIGHT_INDIRECT_NEXT_ITER); } else if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_VOLUME_INDIRECT)) { diff --git a/intern/cycles/kernel/split/kernel_split_data_types.h b/intern/cycles/kernel/split/kernel_split_data_types.h index bb1aca2acbf..4bb2f0d3d80 100644 --- a/intern/cycles/kernel/split/kernel_split_data_types.h +++ b/intern/cycles/kernel/split/kernel_split_data_types.h @@ -95,6 +95,10 @@ typedef ccl_global struct SplitBranchedState { VolumeStack volume_stack[VOLUME_STACK_SIZE]; # endif /* __VOLUME__ */ #endif /*__SUBSURFACE__ */ + + int shared_sample_count; /* number of branched samples shared with other threads */ + int original_ray; /* index of original ray when sharing branched samples */ + bool waiting_on_shared_samples; } SplitBranchedState; #define SPLIT_DATA_BRANCHED_ENTRIES \ @@ -137,6 +141,25 @@ typedef ccl_global struct SplitBranchedState { SPLIT_DATA_BRANCHED_ENTRIES \ SPLIT_DATA_DEBUG_ENTRIES \ +/* entries to be copied to inactive rays when sharing branched samples (TODO: which are actually needed?) */ +#define SPLIT_DATA_ENTRIES_BRANCHED_SHARED \ + SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \ + SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \ + SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \ + SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \ + SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \ + SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \ + SPLIT_DATA_ENTRY(ccl_global Intersection, isect, 1) \ + SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ + SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ + SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ + SPLIT_DATA_ENTRY(ShaderData, sd, 1) \ + SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 1) \ + SPLIT_DATA_SUBSURFACE_ENTRIES \ + SPLIT_DATA_VOLUME_ENTRIES \ + SPLIT_DATA_BRANCHED_ENTRIES \ + SPLIT_DATA_DEBUG_ENTRIES \ + /* struct that holds pointers to data in the shared state buffer */ typedef struct SplitData { #define SPLIT_DATA_ENTRY(type, name, num) type *name; diff --git a/intern/cycles/kernel/split/kernel_subsurface_scatter.h b/intern/cycles/kernel/split/kernel_subsurface_scatter.h index 1dffe1b179e..4998714f28c 100644 --- a/intern/cycles/kernel/split/kernel_subsurface_scatter.h +++ b/intern/cycles/kernel/split/kernel_subsurface_scatter.h @@ -169,6 +169,7 @@ ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_it ray_index, num_samples_inv, bssrdf_sd, + false, false)) { branched_state->ss_next_closure = i; @@ -187,6 +188,13 @@ ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_it branched_state->ss_next_sample = 0; } + branched_state->ss_next_closure = sd->num_closure; + + branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); + if(branched_state->waiting_on_shared_samples) { + return true; + } + kernel_split_branched_path_indirect_loop_end(kg, ray_index); return false; diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp index 08909943c49..ae462a1084a 100644 --- a/intern/cycles/render/session.cpp +++ b/intern/cycles/render/session.cpp @@ -722,6 +722,7 @@ DeviceRequestedFeatures Session::get_requested_device_features() requested_features.use_baking = bake_manager->get_baking(); requested_features.use_integrator_branched = (scene->integrator->method == Integrator::BRANCHED_PATH); requested_features.use_transparent &= scene->integrator->transparent_shadows; + requested_features.use_denoising = params.use_denoising; return requested_features; } diff --git a/intern/cycles/util/util_atomic.h b/intern/cycles/util/util_atomic.h index 6c52117ef9a..643af87a65f 100644 --- a/intern/cycles/util/util_atomic.h +++ b/intern/cycles/util/util_atomic.h @@ -35,6 +35,7 @@ ATOMIC_INLINE void atomic_update_max_z(size_t *maximum_value, size_t value) #define atomic_add_and_fetch_float(p, x) atomic_add_and_fetch_fl((p), (x)) #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1) +#define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_add_uint32((p), -1) #define CCL_LOCAL_MEM_FENCE 0 #define ccl_barrier(flags) (void)0 @@ -68,6 +69,7 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so #define atomic_fetch_and_add_uint32(p, x) atomic_add((p), (x)) #define atomic_fetch_and_inc_uint32(p) atomic_inc((p)) +#define atomic_fetch_and_dec_uint32(p) atomic_dec((p)) #define CCL_LOCAL_MEM_FENCE CLK_LOCAL_MEM_FENCE #define ccl_barrier(flags) barrier(flags) @@ -79,7 +81,9 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so #define atomic_add_and_fetch_float(p, x) (atomicAdd((float*)(p), (float)(x)) + (float)(x)) #define atomic_fetch_and_add_uint32(p, x) atomicAdd((unsigned int*)(p), (unsigned int)(x)) +#define atomic_fetch_and_sub_uint32(p, x) atomicSub((unsigned int*)(p), (unsigned int)(x)) #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1) +#define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_sub_uint32((p), 1) #define CCL_LOCAL_MEM_FENCE #define ccl_barrier(flags) __syncthreads() diff --git a/release/datafiles/blender_icons.svg b/release/datafiles/blender_icons.svg index d88788fa904..a9c0fd431eb 100644 --- a/release/datafiles/blender_icons.svg +++ b/release/datafiles/blender_icons.svg @@ -31358,6 +31358,136 @@ fx="-0.78262758" fy="294.63174" r="6.6750002" /> + <radialGradient + inkscape:collect="always" + xlink:href="#linearGradient33897" + id="radialGradient29130" + gradientUnits="userSpaceOnUse" + gradientTransform="matrix(1.6249996,2.7477764e-7,-3.1704883e-7,1.874986,-24.234082,-761.21063)" + cx="39.528847" + cy="871.2453" + fx="39.528847" + fy="871.2453" + r="2.0000005" /> + <linearGradient + inkscape:collect="always" + id="linearGradient33897"> + <stop + style="stop-color:#ffffff;stop-opacity:1;" + offset="0" + id="stop33893" /> + <stop + style="stop-color:#ffffff;stop-opacity:0;" + offset="1" + id="stop33895" /> + </linearGradient> + <linearGradient + inkscape:collect="always" + xlink:href="#linearGradient319-367" + id="linearGradient63713" + gradientUnits="userSpaceOnUse" + gradientTransform="matrix(0.8167109,0,0,0.8433415,239.34332,-149.78578)" + x1="104.90227" + y1="53.227627" + x2="114.94328" + y2="60.73848" /> + <linearGradient + id="linearGradient319-367"> + <stop + style="stop-color:#ffffff;stop-opacity:1;" + offset="0" + id="stop320-53" /> + <stop + style="stop-color:#ffffff;stop-opacity:0;" + offset="1" + id="stop321-562" /> + </linearGradient> + <linearGradient + inkscape:collect="always" + xlink:href="#linearGradient319-367" + id="linearGradient63715" + gradientUnits="userSpaceOnUse" + gradientTransform="translate(207,-246.99988)" + x1="-56.65625" + y1="342.03125" + x2="-53.1875" + y2="342.0625" /> + <linearGradient + id="linearGradient3043"> + <stop + style="stop-color:#ffffff;stop-opacity:1;" + offset="0" + id="stop3045" /> + <stop + style="stop-color:#ffffff;stop-opacity:0;" + offset="1" + id="stop3047" /> + </linearGradient> + <linearGradient + inkscape:collect="always" + xlink:href="#linearGradient319-367" + id="linearGradient63886" + x1="149.53125" + y1="95.781372" + x2="149.40625" + y2="103.12512" + gradientUnits="userSpaceOnUse" /> + <linearGradient + id="linearGradient3050"> + <stop + style="stop-color:#ffffff;stop-opacity:1;" + offset="0" + id="stop3052" /> + <stop + style="stop-color:#ffffff;stop-opacity:0;" + offset="1" + id="stop3054" /> + </linearGradient> + <linearGradient + inkscape:collect="always" + xlink:href="#linearGradient63892" + id="linearGradient63894" + x1="127.85783" + y1="115.03898" + x2="137.88899" + y2="121.44501" + gradientUnits="userSpaceOnUse" /> + <linearGradient + inkscape:collect="always" + id="linearGradient63892"> + <stop + style="stop-color:#ba5d00;stop-opacity:1" + offset="0" + id="stop63888" /> + <stop + id="stop63896" + offset="0.5" + style="stop-color:#fa9a3a;stop-opacity:1" /> + <stop + style="stop-color:#ffffff;stop-opacity:1" + offset="1" + id="stop63890" /> + </linearGradient> + <linearGradient + inkscape:collect="always" + xlink:href="#linearGradient319-367" + id="linearGradient63719" + gradientUnits="userSpaceOnUse" + x1="132" + y1="117.26753" + x2="142.72656" + y2="127.72736" /> + <linearGradient + id="linearGradient3062"> + <stop + style="stop-color:#ffffff;stop-opacity:1;" + offset="0" + id="stop3064" /> + <stop + style="stop-color:#ffffff;stop-opacity:0;" + offset="1" + id="stop3066" /> + </linearGradient> </defs> <sodipodi:namedview id="base" @@ -92721,6 +92851,406 @@ style="display:inline;overflow:visible;visibility:visible;fill:none;stroke:#a8df84;stroke-width:2.93474906;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;enable-background:accumulate" id="path39834-2" inkscape:connector-curvature="0" /> + <g + id="g1418" + transform="matrix(1.6674711,0,0,1.6674711,-416.35106,-776.00982)" + style="display:inline;enable-background:new"> + <path + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1374" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="4.2499995" + sodipodi:ry="4.2499995" + sodipodi:start="6.1086524" + sodipodi:end="7.5049158" + sodipodi:open="true" + d="m 45.185432,872.62412 c 0.358149,2.03116 -0.793733,4.02628 -2.731847,4.73169" /> + <path + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1376" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="7" + sodipodi:ry="7" + sodipodi:start="4.4505896" + sodipodi:end="4.9741884" + sodipodi:open="true" + d="m 39.188267,866.60064 c 1.186888,-0.31802 2.436579,-0.31802 3.623467,0" /> + <path + d="m 47.76148,875.17385 c -0.318025,1.18689 -0.942871,2.26916 -1.811733,3.13802" + sodipodi:open="true" + sodipodi:end="0.78539816" + sodipodi:start="0.26179939" + sodipodi:ry="6.9999995" + sodipodi:rx="6.9999995" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1378" + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + d="m 39.546414,877.35581 c -1.938113,-0.70541 -3.089995,-2.70053 -2.731846,-4.73169" + sodipodi:open="true" + sodipodi:end="3.3161256" + sodipodi:start="1.9198622" + sodipodi:ry="4.2499995" + sodipodi:rx="4.2499995" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1380" + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1382" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="6.9999995" + sodipodi:ry="6.9999995" + sodipodi:start="2.3561945" + sodipodi:end="2.8797933" + sodipodi:open="true" + d="m 36.050253,878.31187 c -0.868862,-0.86886 -1.493708,-1.95113 -1.811733,-3.13802" /> + <rect + ry="2.75" + rx="2.75" + y="870.61212" + x="38.25" + height="5.5" + width="5.5000005" + id="rect1384" + style="fill:#4b2f1e;fill-opacity:1;fill-rule:nonzero;stroke:none" /> + <rect + style="fill:#000000;fill-opacity:0;fill-rule:nonzero;stroke:none" + id="rect1386" + width="20" + height="20" + x="31" + y="862.36212" + rx="0" + ry="0" /> + <rect + style="fill:#ff7d1f;fill-opacity:1;fill-rule:nonzero;stroke:none" + id="rect1388" + width="4.0000005" + height="4" + x="39" + y="871.36212" + rx="2" + ry="2" /> + <rect + ry="1.5" + rx="1.5000004" + y="871.86212" + x="39.5" + height="3" + width="3.0000007" + id="rect1390" + style="fill:none;stroke:url(#radialGradient29130);stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1392" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="4.2499995" + sodipodi:ry="4.2499995" + sodipodi:start="1.9198622" + sodipodi:end="3.3161256" + sodipodi:open="true" + d="m 39.546414,877.35581 c -1.938113,-0.70541 -3.089995,-2.70053 -2.731846,-4.73169" /> + <path + d="m 36.050253,878.31187 c -0.868862,-0.86886 -1.493708,-1.95113 -1.811733,-3.13802" + sodipodi:open="true" + sodipodi:end="2.8797933" + sodipodi:start="2.3561945" + sodipodi:ry="6.9999995" + sodipodi:rx="6.9999995" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1394" + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + d="m 38.875,877.04273 c -0.321653,-0.18571 -0.617575,-0.41278 -0.880204,-0.6754" + sodipodi:open="true" + sodipodi:end="2.3561945" + sodipodi:start="2.0943951" + sodipodi:ry="4.2499995" + sodipodi:rx="4.2499995" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1396" + style="opacity:0.7;fill:none;stroke:#ff8919;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + style="opacity:0.7;fill:none;stroke:#ff8919;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1398" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="6.9999995" + sodipodi:ry="6.9999995" + sodipodi:start="2.3561945" + sodipodi:end="2.6179939" + sodipodi:open="true" + d="m 36.050253,878.31187 c -0.432565,-0.43257 -0.806561,-0.91997 -1.11243,-1.44975" /> + <path + style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1400" + sodipodi:type="arc" + sodipodi:cx="-41" + sodipodi:cy="-873.36212" + sodipodi:rx="4.2500005" + sodipodi:ry="4.2500005" + sodipodi:start="0.87266463" + sodipodi:end="2.268928" + sodipodi:open="true" + d="m -38.268152,-870.10643 c -1.579966,1.32575 -3.88373,1.32575 -5.463696,0" + inkscape:transform-center-x="-2.2499995" + inkscape:transform-center-y="-2.25" + transform="scale(-1,-1)" /> + <path + d="m -38.268152,-870.10643 c -1.579966,1.32575 -3.88373,1.32575 -5.463696,0" + sodipodi:open="true" + sodipodi:end="2.268928" + sodipodi:start="0.87266463" + sodipodi:ry="4.2500005" + sodipodi:rx="4.2500005" + sodipodi:cy="-873.36212" + sodipodi:cx="-41" + sodipodi:type="arc" + id="path1402" + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + inkscape:transform-center-x="-2.2499995" + inkscape:transform-center-y="-2.25" + transform="scale(-1,-1)" /> + <path + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1404" + sodipodi:type="arc" + sodipodi:cx="-41" + sodipodi:cy="-873.36212" + sodipodi:rx="7.0000005" + sodipodi:ry="7.0000005" + sodipodi:start="1.3089969" + sodipodi:end="1.8325957" + sodipodi:open="true" + d="m -39.188266,-866.60064 c -1.186888,0.31803 -2.436579,0.31803 -3.623467,0" + inkscape:transform-center-y="-3.7500001" + transform="scale(-1,-1)" + inkscape:transform-center-x="-3.7499995" /> + <path + style="opacity:0.3;fill:none;stroke:#ffe4cb;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1406" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="4.25" + sodipodi:ry="4.25" + sodipodi:start="4.0142573" + sodipodi:end="4.9741884" + sodipodi:open="true" + d="m 38.268153,870.10643 c 1.062216,-0.8913 2.492452,-1.20838 3.831828,-0.84949" /> + <path + d="m 39.188267,866.60064 c 1.186888,-0.31802 2.436579,-0.31802 3.623467,0" + sodipodi:open="true" + sodipodi:end="4.9741884" + sodipodi:start="4.4505896" + sodipodi:ry="7" + sodipodi:rx="7" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1408" + style="opacity:0.3;fill:none;stroke:#ffe4cb;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + style="opacity:0.6;fill:none;stroke:#ffffff;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1410" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="6.9999995" + sodipodi:ry="6.9999995" + sodipodi:start="4.4505896" + sodipodi:end="4.712389" + sodipodi:open="true" + d="M 39.188267,866.60064 C 39.779161,866.44231 40.388261,866.36212 41,866.36212" /> + <path + d="m 38.268153,870.10643 c 0.475408,-0.39891 1.032412,-0.68887 1.631866,-0.84949" + sodipodi:open="true" + sodipodi:end="4.4505896" + sodipodi:start="4.0142573" + sodipodi:ry="4.25" + sodipodi:rx="4.25" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1412" + style="opacity:0.6;fill:none;stroke:#ffffff;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" /> + <path + sodipodi:open="true" + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="path1414" + sodipodi:type="arc" + sodipodi:cx="41" + sodipodi:cy="873.36212" + sodipodi:rx="7" + sodipodi:ry="7" + sodipodi:start="0.26179939" + sodipodi:end="0.78539816" + d="m 47.761481,875.17385 c -0.318026,1.18689 -0.942871,2.26916 -1.811734,3.13802" /> + <path + d="m 45.185432,872.62412 c 0.358149,2.03116 -0.793733,4.02628 -2.731847,4.73169" + sodipodi:end="7.5049158" + sodipodi:start="6.1086524" + sodipodi:ry="4.2499995" + sodipodi:rx="4.2499995" + sodipodi:cy="873.36212" + sodipodi:cx="41" + sodipodi:type="arc" + id="path1416" + style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + sodipodi:open="true" /> + </g> + <g + transform="matrix(1.6674711,0,0,1.6674711,-138.14039,-258.06137)" + style="display:inline;enable-background:new" + id="g63699"> + <path + inkscape:connector-curvature="0" + style="opacity:0.3;fill:none;stroke:#422200;stroke-width:0.80000001;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + d="m -102.5625,554.75 c -0.0429,0.005 -0.0849,0.0154 -0.125,0.0312 l -4.5,1.75 c -0.1919,0.0756 -0.31653,0.26258 -0.3125,0.4688 v 6.5 c 0.003,0.18741 0.11203,0.35691 0.28125,0.4375 l 4.5,2.25 c 0.13787,0.0682 0.29963,0.0682 0.4375,0 l 4.499997,-2.25 c 0.169224,-0.0806 0.278188,-0.25009 0.28125,-0.4375 V 557 c 0.004,-0.2062 -0.120622,-0.39315 -0.3125,-0.46875 l -4.499997,-1.75 c -0.0792,-0.0318 -0.16537,-0.0426 -0.25,-0.0312 z" + id="path63697" /> + <g + id="g63687" + transform="translate(-252,462.99988)"> + <g + transform="translate(-179,199.50012)" + id="g63679"> + <path + inkscape:connector-curvature="0" + id="path63667" + d="m 328.5,-107.25 -4.5,1.75 v 6.5 l 4.5,2.25 4.5,-2.25 v -6.5 z" + style="fill:#422200;fill-opacity:1;fill-rule:evenodd;stroke:#422200;stroke-width:1.79999995;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + sodipodi:nodetypes="ccccccc" /> + <g + transform="translate(179,-179)" + id="g63671"> + <path + inkscape:connector-curvature="0" + inkscape:export-ydpi="90" + inkscape:export-xdpi="90" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png" + style="fill:#915515;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + d="m 154,80 v -6.5 l -4.5,-1.75 v 10.5 z" + id="path63669" + sodipodi:nodetypes="ccccc" /> + </g> + <path + inkscape:connector-curvature="0" + sodipodi:nodetypes="ccccccc" + style="fill:#efa351;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + d="M 324,-99.00012 V -105.5 l 4.5,-1.75 0.5,0.25 v 10 l -0.5,0.25 z" + id="path63673" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png" + inkscape:export-xdpi="90" + inkscape:export-ydpi="90" /> + <path + inkscape:connector-curvature="0" + style="fill:none;stroke:url(#linearGradient63713);stroke-width:0.99999982px;stroke-linecap:butt;stroke-linejoin:miter;stroke-opacity:1" + d="m 332.5,-105.5 v 6.25 l -4,2 -4,-2.00012 V -105.5" + id="path63675" + sodipodi:nodetypes="ccccc" /> + <path + inkscape:connector-curvature="0" + inkscape:export-ydpi="90" + inkscape:export-xdpi="90" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png" + sodipodi:nodetypes="ccccc" + style="fill:#f5ca9b;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + d="m 324,-105.5 4.5,-1.75 4.5,1.75 -4.5,2 z" + id="path63677" /> + </g> + <path + inkscape:connector-curvature="0" + sodipodi:nodetypes="ccc" + style="opacity:0.95999995;fill:none;stroke:url(#linearGradient63715);stroke-width:1px;stroke-linecap:round;stroke-linejoin:round;stroke-opacity:1" + d="m 145.5,94.25012 c 0,0 4,1.75 4,1.75 l 4,-1.75" + id="path63681" /> + <rect + style="opacity:0.25;fill:url(#linearGradient63886);fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + id="rect63683" + width="1" + height="6.7500019" + x="149" + y="96.000122" /> + </g> + <rect + style="opacity:0;fill:#b3b3b3;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:2.79999995;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate" + id="rect63665" + width="16" + height="16" + x="-113" + y="554" /> + <g + style="display:inline;enable-background:new" + id="g63695" + transform="matrix(0.7882544,0,0,0.7883038,-210.45268,388.9974)" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\BLENDER ICONSET\blender-cvs-windows\.blender\.blender\icons\jendrzych's iconset.png" + inkscape:export-xdpi="90" + inkscape:export-ydpi="90"> + <circle + sodipodi:ry="8" + sodipodi:rx="8" + sodipodi:cy="118" + sodipodi:cx="132" + r="8" + cy="118" + cx="132" + inkscape:export-ydpi="90" + inkscape:export-xdpi="90" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png" + id="circle63689" + style="fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:#422200;stroke-width:1.97436094;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + transform="matrix(0.6425292,0,0,0.642531,44.523834,146.81699)" + d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" /> + <circle + sodipodi:ry="8" + sodipodi:rx="8" + sodipodi:cy="118" + sodipodi:cx="132" + r="8" + cy="118" + cx="132" + style="opacity:0.8;fill:url(#linearGradient63894);fill-opacity:1;fill-rule:nonzero;stroke:none;display:inline" + id="circle63691" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png" + inkscape:export-xdpi="90" + inkscape:export-ydpi="90" + transform="matrix(-0.5858806,-0.06590218,0.06677852,-0.5812167,198.80048,299.96262)" + d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" /> + <circle + sodipodi:ry="8" + sodipodi:rx="8" + sodipodi:cy="118" + sodipodi:cx="132" + r="8" + cy="118" + cx="132" + transform="matrix(0.4991181,0,0,0.4991107,63.460522,163.7471)" + style="opacity:0.5;fill:none;stroke:url(#linearGradient63719);stroke-width:2.54167628;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" + id="circle63693" + inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png" + inkscape:export-xdpi="90" + inkscape:export-ydpi="90" + d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" /> + </g> + </g> </g> <path id="path39836-9" diff --git a/release/datafiles/blender_icons16/icon16_outliner_ob_force_field.dat b/release/datafiles/blender_icons16/icon16_outliner_ob_force_field.dat Binary files differnew file mode 100644 index 00000000000..59817bce4de --- /dev/null +++ b/release/datafiles/blender_icons16/icon16_outliner_ob_force_field.dat diff --git a/release/datafiles/blender_icons16/icon16_outliner_ob_group_instance.dat b/release/datafiles/blender_icons16/icon16_outliner_ob_group_instance.dat Binary files differnew file mode 100644 index 00000000000..072c049c941 --- /dev/null +++ b/release/datafiles/blender_icons16/icon16_outliner_ob_group_instance.dat diff --git a/release/datafiles/blender_icons32/icon32_outliner_ob_force_field.dat b/release/datafiles/blender_icons32/icon32_outliner_ob_force_field.dat Binary files differnew file mode 100644 index 00000000000..bd57178b377 --- /dev/null +++ b/release/datafiles/blender_icons32/icon32_outliner_ob_force_field.dat diff --git a/release/datafiles/blender_icons32/icon32_outliner_ob_group_instance.dat b/release/datafiles/blender_icons32/icon32_outliner_ob_group_instance.dat Binary files differnew file mode 100644 index 00000000000..868f9e4fbcc --- /dev/null +++ b/release/datafiles/blender_icons32/icon32_outliner_ob_group_instance.dat diff --git a/release/scripts/modules/bl_i18n_utils/settings.py b/release/scripts/modules/bl_i18n_utils/settings.py index 07fc9f0c338..774ca5173d5 100644 --- a/release/scripts/modules/bl_i18n_utils/settings.py +++ b/release/scripts/modules/bl_i18n_utils/settings.py @@ -118,7 +118,7 @@ MSG_COMMENT_PREFIX = "#~ " MSG_CONTEXT_PREFIX = "MSGCTXT:" # The default comment prefix used in po's. -PO_COMMENT_PREFIX= "# " +PO_COMMENT_PREFIX = "# " # The comment prefix used to mark sources of msgids, in po's. PO_COMMENT_PREFIX_SOURCE = "#: " @@ -130,7 +130,7 @@ PO_COMMENT_PREFIX_SOURCE_CUSTOM = "#. :src: " PO_COMMENT_PREFIX_GENERATED = "#. " # The comment prefix used to comment entries in po's. -PO_COMMENT_PREFIX_MSG= "#~ " +PO_COMMENT_PREFIX_MSG = "#~ " # The comment prefix used to mark fuzzy msgids, in po's. PO_COMMENT_FUZZY = "#, fuzzy" diff --git a/release/scripts/startup/bl_ui/properties_material.py b/release/scripts/startup/bl_ui/properties_material.py index efdd22c5d0c..2feb60f3bfa 100644 --- a/release/scripts/startup/bl_ui/properties_material.py +++ b/release/scripts/startup/bl_ui/properties_material.py @@ -124,9 +124,10 @@ class MATERIAL_PT_context_material(MaterialButtonsPanel, Panel): ob = context.object slot = context.material_slot space = context.space_data - is_sortable = (len(ob.material_slots) > 1) if ob: + is_sortable = (len(ob.material_slots) > 1) + rows = 1 if is_sortable: rows = 4 diff --git a/release/scripts/startup/bl_ui/space_sequencer.py b/release/scripts/startup/bl_ui/space_sequencer.py index 6737b8e1089..79bb10cefeb 100644 --- a/release/scripts/startup/bl_ui/space_sequencer.py +++ b/release/scripts/startup/bl_ui/space_sequencer.py @@ -338,21 +338,21 @@ class SEQUENCER_MT_add(Menu): if len(bpy.data.scenes) > 10: layout.operator_context = 'INVOKE_DEFAULT' - layout.operator("sequencer.scene_strip_add", text="Scene") + layout.operator("sequencer.scene_strip_add", text="Scene...") else: - layout.operator_menu_enum("sequencer.scene_strip_add", "scene", text="Scene...") + layout.operator_menu_enum("sequencer.scene_strip_add", "scene", text="Scene") if len(bpy.data.movieclips) > 10: layout.operator_context = 'INVOKE_DEFAULT' - layout.operator("sequencer.movieclip_strip_add", text="Clips") + layout.operator("sequencer.movieclip_strip_add", text="Clips...") else: - layout.operator_menu_enum("sequencer.movieclip_strip_add", "clip", text="Clip...") + layout.operator_menu_enum("sequencer.movieclip_strip_add", "clip", text="Clip") if len(bpy.data.masks) > 10: layout.operator_context = 'INVOKE_DEFAULT' - layout.operator("sequencer.mask_strip_add", text="Masks") + layout.operator("sequencer.mask_strip_add", text="Masks...") else: - layout.operator_menu_enum("sequencer.mask_strip_add", "mask", text="Mask...") + layout.operator_menu_enum("sequencer.mask_strip_add", "mask", text="Mask") layout.operator("sequencer.movie_strip_add", text="Movie") layout.operator("sequencer.image_strip_add", text="Image") diff --git a/release/scripts/startup/bl_ui/space_time.py b/release/scripts/startup/bl_ui/space_time.py index b9a25cd72a0..9930992e9bb 100644 --- a/release/scripts/startup/bl_ui/space_time.py +++ b/release/scripts/startup/bl_ui/space_time.py @@ -250,7 +250,7 @@ def marker_menu_generic(layout): layout.operator_context = 'INVOKE_DEFAULT' layout.operator("marker.make_links_scene", text="Duplicate Marker to Scene...", icon='OUTLINER_OB_EMPTY') else: - layout.operator_menu_enum("marker.make_links_scene", "scene", text="Duplicate Marker to Scene...") + layout.operator_menu_enum("marker.make_links_scene", "scene", text="Duplicate Marker to Scene") layout.operator("marker.delete", text="Delete Marker") diff --git a/release/scripts/startup/bl_ui/space_userpref.py b/release/scripts/startup/bl_ui/space_userpref.py index 0250f4a8ce8..13a2cb97b2d 100644 --- a/release/scripts/startup/bl_ui/space_userpref.py +++ b/release/scripts/startup/bl_ui/space_userpref.py @@ -129,7 +129,7 @@ class USERPREF_MT_app_templates(Menu): text=bpy.path.display_name(d), ) props.use_splash = True - props.app_template = d; + props.app_template = d if use_install: layout.separator() @@ -561,7 +561,7 @@ class USERPREF_PT_system(Panel): # 3. Column column = split.column() - column.label(text="Solid OpenGL lights:") + column.label(text="Solid OpenGL Lights:") split = column.split(percentage=0.1) split.label() diff --git a/release/scripts/startup/bl_ui/space_view3d.py b/release/scripts/startup/bl_ui/space_view3d.py index a6bd29d19cf..0e1a142a41f 100644 --- a/release/scripts/startup/bl_ui/space_view3d.py +++ b/release/scripts/startup/bl_ui/space_view3d.py @@ -1266,14 +1266,14 @@ class INFO_MT_add(Menu): layout.menu("INFO_MT_lightprobe_add") layout.separator() - layout.operator_menu_enum("object.effector_add", "type", text="Force Field", icon='OUTLINER_OB_EMPTY') + layout.operator_menu_enum("object.effector_add", "type", text="Force Field", icon='OUTLINER_OB_FORCE_FIELD') layout.separator() if len(bpy.data.groups) > 10: layout.operator_context = 'INVOKE_REGION_WIN' - layout.operator("object.group_instance_add", text="Group Instance...", icon='OUTLINER_OB_EMPTY') + layout.operator("object.group_instance_add", text="Group Instance...", icon='OUTLINER_OB_GROUP_INSTANCE') else: - layout.operator_menu_enum("object.group_instance_add", "group", text="Group Instance", icon='OUTLINER_OB_EMPTY') + layout.operator_menu_enum("object.group_instance_add", "group", text="Group Instance", icon='OUTLINER_OB_GROUP_INSTANCE') # ********** Object menu ********** @@ -1654,7 +1654,7 @@ class VIEW3D_MT_make_links(Menu): layout.operator("object.make_links_scene", text="Objects to Scene...", icon='OUTLINER_OB_EMPTY') else: layout.operator_context = 'EXEC_REGION_WIN' - layout.operator_menu_enum("object.make_links_scene", "scene", text="Objects to Scene...") + layout.operator_menu_enum("object.make_links_scene", "scene", text="Objects to Scene") layout.operator_context = operator_context_default layout.operator_enum("object.make_links_data", "type") # inline diff --git a/source/blender/blenkernel/BKE_animsys.h b/source/blender/blenkernel/BKE_animsys.h index a67e903877a..420ae4bb12d 100644 --- a/source/blender/blenkernel/BKE_animsys.h +++ b/source/blender/blenkernel/BKE_animsys.h @@ -89,12 +89,6 @@ typedef enum eAnimData_MergeCopy_Modes { void BKE_animdata_merge_copy(struct ID *dst_id, struct ID *src_id, eAnimData_MergeCopy_Modes action_mode, bool fix_drivers); -/* Make Local */ -void BKE_animdata_make_local(struct AnimData *adt); - -/* Re-Assign ID's */ -void BKE_animdata_relink(struct AnimData *adt); - /* ************************************* */ /* KeyingSets API */ diff --git a/source/blender/blenkernel/intern/anim_sys.c b/source/blender/blenkernel/intern/anim_sys.c index e455b02fa6e..6475e5c8acd 100644 --- a/source/blender/blenkernel/intern/anim_sys.c +++ b/source/blender/blenkernel/intern/anim_sys.c @@ -399,73 +399,6 @@ void BKE_animdata_merge_copy(ID *dst_id, ID *src_id, eAnimData_MergeCopy_Modes a } } -/* Make Local -------------------------------------------- */ - -static void make_local_strips(ListBase *strips) -{ - NlaStrip *strip; - - for (strip = strips->first; strip; strip = strip->next) { - if (strip->act) BKE_action_make_local(G.main, strip->act, false); - if (strip->remap && strip->remap->target) BKE_action_make_local(G.main, strip->remap->target, false); - - make_local_strips(&strip->strips); - } -} - -/* Use local copy instead of linked copy of various ID-blocks */ -void BKE_animdata_make_local(AnimData *adt) -{ - NlaTrack *nlt; - - /* Actions - Active and Temp */ - if (adt->action) BKE_action_make_local(G.main, adt->action, false); - if (adt->tmpact) BKE_action_make_local(G.main, adt->tmpact, false); - /* Remaps */ - if (adt->remap && adt->remap->target) BKE_action_make_local(G.main, adt->remap->target, false); - - /* Drivers */ - /* TODO: need to remap the ID-targets too? */ - - /* NLA Data */ - for (nlt = adt->nla_tracks.first; nlt; nlt = nlt->next) - make_local_strips(&nlt->strips); -} - - -/* When duplicating data (i.e. objects), drivers referring to the original data will - * get updated to point to the duplicated data (if drivers belong to the new data) - */ -void BKE_animdata_relink(AnimData *adt) -{ - /* sanity check */ - if (adt == NULL) - return; - - /* drivers */ - if (adt->drivers.first) { - FCurve *fcu; - - /* check each driver against all the base paths to see if any should go */ - for (fcu = adt->drivers.first; fcu; fcu = fcu->next) { - ChannelDriver *driver = fcu->driver; - DriverVar *dvar; - - /* driver variables */ - for (dvar = driver->variables.first; dvar; dvar = dvar->next) { - /* only change the used targets, since the others will need fixing manually anyway */ - DRIVER_TARGETS_USED_LOOPER(dvar) - { - if (dtar->id && dtar->id->newid) { - dtar->id = dtar->id->newid; - } - } - DRIVER_TARGETS_LOOPER_END - } - } - } -} - /* Sub-ID Regrouping ------------------------------------------- */ /** diff --git a/source/blender/blenkernel/intern/constraint.c b/source/blender/blenkernel/intern/constraint.c index 9bc1ce16284..07a6b304dff 100644 --- a/source/blender/blenkernel/intern/constraint.c +++ b/source/blender/blenkernel/intern/constraint.c @@ -790,7 +790,7 @@ static void default_get_tarmat(bConstraint *con, bConstraintOb *UNUSED(cob), bCo ct = ctn; \ } \ } (void)0 - + /* --------- ChildOf Constraint ------------ */ static void childof_new_data(void *cdata) diff --git a/source/blender/blenkernel/intern/effect.c b/source/blender/blenkernel/intern/effect.c index 4b5548c48d4..2213094cd0b 100644 --- a/source/blender/blenkernel/intern/effect.c +++ b/source/blender/blenkernel/intern/effect.c @@ -986,19 +986,19 @@ static void do_physical_effector(EffectorCache *eff, EffectorData *efd, Effected */ void pdDoEffectors(ListBase *effectors, ListBase *colliders, EffectorWeights *weights, EffectedPoint *point, float *force, float *impulse) { -/* - * Modifies the force on a particle according to its - * relation with the effector object - * Different kind of effectors include: - * Forcefields: Gravity-like attractor - * (force power is related to the inverse of distance to the power of a falloff value) - * Vortex fields: swirling effectors - * (particles rotate around Z-axis of the object. otherwise, same relation as) - * (Forcefields, but this is not done through a force/acceleration) - * Guide: particles on a path - * (particles are guided along a curve bezier or old nurbs) - * (is independent of other effectors) - */ + /* + * Modifies the force on a particle according to its + * relation with the effector object + * Different kind of effectors include: + * Forcefields: Gravity-like attractor + * (force power is related to the inverse of distance to the power of a falloff value) + * Vortex fields: swirling effectors + * (particles rotate around Z-axis of the object. otherwise, same relation as) + * (Forcefields, but this is not done through a force/acceleration) + * Guide: particles on a path + * (particles are guided along a curve bezier or old nurbs) + * (is independent of other effectors) + */ EffectorCache *eff; EffectorData efd; int p=0, tot = 1, step = 1; diff --git a/source/blender/blenkernel/intern/image.c b/source/blender/blenkernel/intern/image.c index 7ee4afd052d..4d3e2bbc3c8 100644 --- a/source/blender/blenkernel/intern/image.c +++ b/source/blender/blenkernel/intern/image.c @@ -2865,7 +2865,7 @@ bool BKE_image_is_stereo(Image *ima) { return BKE_image_is_multiview(ima) && (BLI_findstring(&ima->views, STEREO_LEFT_NAME, offsetof(ImageView, name)) && - BLI_findstring(&ima->views, STEREO_RIGHT_NAME, offsetof(ImageView, name))); + BLI_findstring(&ima->views, STEREO_RIGHT_NAME, offsetof(ImageView, name))); } static void image_init_multilayer_multiview(Image *ima, RenderResult *rr) diff --git a/source/blender/blenkernel/intern/ocean.c b/source/blender/blenkernel/intern/ocean.c index 621ac9c2480..537c8926a5b 100644 --- a/source/blender/blenkernel/intern/ocean.c +++ b/source/blender/blenkernel/intern/ocean.c @@ -334,10 +334,10 @@ void BKE_ocean_eval_uv(struct Ocean *oc, struct OceanResult *ocr, float u, float i1 = i1 % oc->_M; j1 = j1 % oc->_N; - #define BILERP(m) (interpf(interpf(m[i1 * oc->_N + j1], m[i0 * oc->_N + j1], frac_x), \ interpf(m[i1 * oc->_N + j0], m[i0 * oc->_N + j0], frac_x), \ frac_z)) + { if (oc->_do_disp_y) { ocr->disp[1] = BILERP(oc->_disp_y); diff --git a/source/blender/blenkernel/intern/particle.c b/source/blender/blenkernel/intern/particle.c index 2084e7ec434..df48c319064 100644 --- a/source/blender/blenkernel/intern/particle.c +++ b/source/blender/blenkernel/intern/particle.c @@ -592,7 +592,7 @@ void psys_free(Object *ob, ParticleSystem *psys) BLI_bvhtree_free(psys->bvhtree); BLI_kdtree_free(psys->tree); - + if (psys->fluid_springs) MEM_freeN(psys->fluid_springs); diff --git a/source/blender/blenkernel/intern/pbvh.c b/source/blender/blenkernel/intern/pbvh.c index 33901d3c0d0..336d4fae1b8 100644 --- a/source/blender/blenkernel/intern/pbvh.c +++ b/source/blender/blenkernel/intern/pbvh.c @@ -1114,7 +1114,7 @@ static void pbvh_update_draw_buffers(PBVH *bvh, PBVHNode **nodes, int totnode) GPU_pbvh_bmesh_buffers_build(bvh->flags & PBVH_DYNTOPO_SMOOTH_SHADING); break; } - + node->flag &= ~PBVH_RebuildDrawBuffers; } diff --git a/source/blender/blenkernel/intern/seqeffects.c b/source/blender/blenkernel/intern/seqeffects.c index 298671beedb..e435d87024e 100644 --- a/source/blender/blenkernel/intern/seqeffects.c +++ b/source/blender/blenkernel/intern/seqeffects.c @@ -684,7 +684,7 @@ static float invGammaCorrect(float c) else if (i >= RE_GAMMA_TABLE_SIZE) res = powf(c, valid_inv_gamma); else res = inv_gamma_range_table[i] + ((c - color_domain_table[i]) * inv_gamfactor_table[i]); - + return res; } diff --git a/source/blender/blenkernel/intern/sequencer.c b/source/blender/blenkernel/intern/sequencer.c index 2df601ee4f5..0d7a374a446 100644 --- a/source/blender/blenkernel/intern/sequencer.c +++ b/source/blender/blenkernel/intern/sequencer.c @@ -1777,7 +1777,7 @@ static ImBuf *seq_proxy_fetch(const SeqRenderData *context, Sequence *seq, int c if (proxy->anim == NULL) { return NULL; } - + seq_open_anim_file(context->scene, seq, true); sanim = seq->anims.first; @@ -1785,7 +1785,7 @@ static ImBuf *seq_proxy_fetch(const SeqRenderData *context, Sequence *seq, int c return IMB_anim_absolute(proxy->anim, frameno, IMB_TC_NONE, IMB_PROXY_NONE); } - + if (seq_proxy_get_fname(ed, seq, cfra, render_size, name, context->view_id) == 0) { return NULL; } diff --git a/source/blender/blenkernel/intern/softbody.c b/source/blender/blenkernel/intern/softbody.c index 11c692077c0..0b8c11f6dc8 100644 --- a/source/blender/blenkernel/intern/softbody.c +++ b/source/blender/blenkernel/intern/softbody.c @@ -2235,9 +2235,9 @@ static void sb_cf_threads_run(Scene *scene, Object *ob, float forcetime, float t static void softbody_calc_forcesEx(Scene *scene, SceneLayer *sl, Object *ob, float forcetime, float timenow) { -/* rule we never alter free variables :bp->vec bp->pos in here ! - * this will ruin adaptive stepsize AKA heun! (BM) - */ + /* rule we never alter free variables :bp->vec bp->pos in here ! + * this will ruin adaptive stepsize AKA heun! (BM) + */ SoftBody *sb= ob->soft; /* is supposed to be there */ /*BodyPoint *bproot;*/ /* UNUSED */ ListBase *do_effector = NULL; diff --git a/source/blender/blenkernel/intern/tracking_stabilize.c b/source/blender/blenkernel/intern/tracking_stabilize.c index 3dfaa1ed77d..b8dfb217c16 100644 --- a/source/blender/blenkernel/intern/tracking_stabilize.c +++ b/source/blender/blenkernel/intern/tracking_stabilize.c @@ -203,7 +203,7 @@ static float get_animated_scaleinf(StabContext *ctx, int framenr) static void get_animated_target_pos(StabContext *ctx, int framenr, - float target_pos[2]) + float target_pos[2]) { target_pos[0] = fetch_from_fcurve(ctx->target_pos[0], framenr, diff --git a/source/blender/blenlib/BLI_array.h b/source/blender/blenlib/BLI_array.h index c645ff06c00..74f24c808ff 100644 --- a/source/blender/blenlib/BLI_array.h +++ b/source/blender/blenlib/BLI_array.h @@ -137,8 +137,8 @@ void _bli_array_grow_func(void **arr_p, const void *arr_static, #define BLI_array_free(arr) \ if (arr && (char *)arr != _##arr##_static) { \ - BLI_array_fake_user(arr); \ - MEM_freeN(arr); \ + BLI_array_fake_user(arr); \ + MEM_freeN(arr); \ } (void)0 #define BLI_array_pop(arr) ( \ diff --git a/source/blender/blenlib/BLI_kdopbvh.h b/source/blender/blenlib/BLI_kdopbvh.h index ba565fca522..564659ad21e 100644 --- a/source/blender/blenlib/BLI_kdopbvh.h +++ b/source/blender/blenlib/BLI_kdopbvh.h @@ -36,7 +36,7 @@ */ #ifdef __cplusplus -extern "C" { +extern "C" { #endif struct BVHTree; @@ -62,7 +62,7 @@ typedef struct BVHTreeNearest { int index; /* the index of the nearest found (untouched if none is found within a dist radius from the given coordinates) */ float co[3]; /* nearest coordinates (untouched it none is found within a dist radius from the given coordinates) */ float no[3]; /* normal at nearest coordinates (untouched it none is found within a dist radius from the given coordinates) */ - float dist_sq; /* squared distance to search arround */ + float dist_sq; /* squared distance to search around */ int flags; } BVHTreeNearest; diff --git a/source/blender/blenlib/intern/BLI_kdopbvh.c b/source/blender/blenlib/intern/BLI_kdopbvh.c index 857c2a5201c..e5ca53a0193 100644 --- a/source/blender/blenlib/intern/BLI_kdopbvh.c +++ b/source/blender/blenlib/intern/BLI_kdopbvh.c @@ -62,6 +62,13 @@ /* used for iterative_raycast */ // #define USE_SKIP_LINKS +/* Use to print balanced output. */ +// #define USE_PRINT_TREE + +/* Check tree is valid. */ +// #define USE_VERIFY_TREE + + #define MAX_TREETYPE 32 /* Setting zero so we can catch bugs in BLI_task/KDOPBVH. @@ -571,10 +578,12 @@ static void node_join(BVHTree *tree, BVHNode *node) } } -/* +#ifdef USE_PRINT_TREE + +/** * Debug and information functions */ -#if 0 + static void bvhtree_print_tree(BVHTree *tree, BVHNode *node, int depth) { int i; @@ -597,30 +606,29 @@ static void bvhtree_print_tree(BVHTree *tree, BVHNode *node, int depth) static void bvhtree_info(BVHTree *tree) { - printf("BVHTree info\n"); - printf("tree_type = %d, axis = %d, epsilon = %f\n", + printf("BVHTree Info: tree_type = %d, axis = %d, epsilon = %f\n", tree->tree_type, tree->axis, tree->epsilon); printf("nodes = %d, branches = %d, leafs = %d\n", tree->totbranch + tree->totleaf, tree->totbranch, tree->totleaf); - printf("Memory per node = %ldbytes\n", - sizeof(BVHNode) + sizeof(BVHNode *) * tree->tree_type + sizeof(float) * tree->axis); - printf("BV memory = %dbytes\n", - (int)MEM_allocN_len(tree->nodebv)); - - printf("Total memory = %ldbytes\n", sizeof(BVHTree) + - MEM_allocN_len(tree->nodes) + - MEM_allocN_len(tree->nodearray) + - MEM_allocN_len(tree->nodechild) + - MEM_allocN_len(tree->nodebv)); - -// bvhtree_print_tree(tree, tree->nodes[tree->totleaf], 0); + printf("Memory per node = %ubytes\n", + (uint)(sizeof(BVHNode) + sizeof(BVHNode *) * tree->tree_type + sizeof(float) * tree->axis)); + printf("BV memory = %ubytes\n", + (uint)MEM_allocN_len(tree->nodebv)); + + printf("Total memory = %ubytes\n", + (uint)(sizeof(BVHTree) + + MEM_allocN_len(tree->nodes) + + MEM_allocN_len(tree->nodearray) + + MEM_allocN_len(tree->nodechild) + + MEM_allocN_len(tree->nodebv))); + + bvhtree_print_tree(tree, tree->nodes[tree->totleaf], 0); } -#endif +#endif /* USE_PRINT_TREE */ -#if 0 +#ifdef USE_VERIFY_TREE - -static void verify_tree(BVHTree *tree) +static void bvhtree_verify(BVHTree *tree) { int i, j, check = 0; @@ -661,7 +669,7 @@ static void verify_tree(BVHTree *tree) printf("branches: %d, leafs: %d, total: %d\n", tree->totbranch, tree->totleaf, tree->totbranch + tree->totleaf); } -#endif +#endif /* USE_VERIFY_TREE */ /* Helper data and structures to build a min-leaf generalized implicit tree * This code can be easily reduced @@ -907,16 +915,24 @@ static void non_recursive_bvh_div_nodes( /* Loop tree levels (log N) loops */ for (i = 1, depth = 1; i <= num_branches; i = i * tree_type + tree_offset, depth++) { const int first_of_next_level = i * tree_type + tree_offset; - const int end_j = min_ii(first_of_next_level, num_branches + 1); /* index of last branch on this level */ + const int i_stop = min_ii(first_of_next_level, num_branches + 1); /* index of last branch on this level */ /* Loop all branches on this level */ cb_data.first_of_next_level = first_of_next_level; cb_data.i = i; cb_data.depth = depth; - BLI_task_parallel_range( - i, end_j, &cb_data, non_recursive_bvh_div_nodes_task_cb, - num_leafs > KDOPBVH_THREAD_LEAF_THRESHOLD); + if (true) { + BLI_task_parallel_range( + i, i_stop, &cb_data, non_recursive_bvh_div_nodes_task_cb, + num_leafs > KDOPBVH_THREAD_LEAF_THRESHOLD); + } + else { + /* Less hassle for debugging. */ + for (int i_task = i; i_task < i_stop; i_task++) { + non_recursive_bvh_div_nodes_task_cb(&cb_data, i_task); + } + } } } @@ -1050,7 +1066,13 @@ void BLI_bvhtree_balance(BVHTree *tree) build_skip_links(tree, tree->nodes[tree->totleaf], NULL, NULL); #endif - /* bvhtree_info(tree); */ +#ifdef USE_VERIFY_TREE + bvhtree_verify(tree); +#endif + +#ifdef USE_PRINT_TREE + bvhtree_info(tree); +#endif } void BLI_bvhtree_insert(BVHTree *tree, int index, const float co[3], int numpoints) diff --git a/source/blender/blenlib/intern/array_store.c b/source/blender/blenlib/intern/array_store.c index 5d1b2433084..d3a63aceb89 100644 --- a/source/blender/blenlib/intern/array_store.c +++ b/source/blender/blenlib/intern/array_store.c @@ -1560,7 +1560,7 @@ BArrayState *BLI_array_store_state_add( const void *data, const size_t data_len, const BArrayState *state_reference) { - /* ensure we're aligned to the stride */ + /* ensure we're aligned to the stride */ BLI_assert((data_len % bs->info.chunk_stride) == 0); #ifdef USE_PARANOID_CHECKS diff --git a/source/blender/blenlib/intern/storage.c b/source/blender/blenlib/intern/storage.c index 48280c14d7d..b819c513fbd 100644 --- a/source/blender/blenlib/intern/storage.c +++ b/source/blender/blenlib/intern/storage.c @@ -422,13 +422,13 @@ void BLI_file_free_lines(LinkNode *lines) bool BLI_file_older(const char *file1, const char *file2) { #ifdef WIN32 - struct _stat st1, st2; + struct _stat st1, st2; UTF16_ENCODE(file1); UTF16_ENCODE(file2); if (_wstat(file1_16, &st1)) return false; - if (_wstat(file2_16, &st2)) return false; + if (_wstat(file2_16, &st2)) return false; UTF16_UN_ENCODE(file2); UTF16_UN_ENCODE(file1); diff --git a/source/blender/blenlib/intern/task.c b/source/blender/blenlib/intern/task.c index a1eae8f1955..e050f3148b8 100644 --- a/source/blender/blenlib/intern/task.c +++ b/source/blender/blenlib/intern/task.c @@ -1127,7 +1127,7 @@ static void task_parallel_range_ex( atomic_fetch_and_add_uint32((uint32_t *)(&state.iter), 0); if (use_userdata_chunk) { - userdata_chunk_array = MALLOCA(userdata_chunk_size * num_tasks); + userdata_chunk_array = MALLOCA(userdata_chunk_size * num_tasks); } for (i = 0; i < num_tasks; i++) { diff --git a/source/blender/blenloader/intern/readfile.c b/source/blender/blenloader/intern/readfile.c index 884bb363333..682ca637fcb 100644 --- a/source/blender/blenloader/intern/readfile.c +++ b/source/blender/blenloader/intern/readfile.c @@ -10374,7 +10374,7 @@ void BLO_library_link_copypaste(Main *mainl, BlendHandle *bh) static ID *link_named_part_ex( Main *mainl, FileData *fd, const short idcode, const char *name, const short flag, - Scene *scene, SceneLayer *sl, const bool use_placeholders, const bool force_indirect) + Scene *scene, SceneLayer *sl, const bool use_placeholders, const bool force_indirect) { ID *id = link_named_part(mainl, fd, idcode, name, use_placeholders, force_indirect); diff --git a/source/blender/depsgraph/DEG_depsgraph_build.h b/source/blender/depsgraph/DEG_depsgraph_build.h index d861afd1fa7..8c3ddec40a4 100644 --- a/source/blender/depsgraph/DEG_depsgraph_build.h +++ b/source/blender/depsgraph/DEG_depsgraph_build.h @@ -153,7 +153,8 @@ void DEG_add_object_cache_relation(struct DepsNodeHandle *handle, eDepsObjectComponentType component, const char *description); -/* TODO(sergey): Remove once all geometry update is granular. */ + +struct Depsgraph *DEG_get_graph_from_handle(struct DepsNodeHandle *handle); void DEG_add_special_eval_flag(struct Depsgraph *graph, struct ID *id, short flag); /* Utility functions for physics modifiers */ diff --git a/source/blender/depsgraph/intern/builder/deg_builder_relations.cc b/source/blender/depsgraph/intern/builder/deg_builder_relations.cc index 9321bd0bc1f..9d1c5816f7c 100644 --- a/source/blender/depsgraph/intern/builder/deg_builder_relations.cc +++ b/source/blender/depsgraph/intern/builder/deg_builder_relations.cc @@ -372,6 +372,11 @@ void DepsgraphRelationBuilder::add_forcefield_relations(const OperationKey &key, pdEndEffectors(&effectors); } +Depsgraph *DepsgraphRelationBuilder::getGraph() +{ + return m_graph; +} + /* **** Functions to build relations between entities **** */ void DepsgraphRelationBuilder::begin_build(Main *bmain) diff --git a/source/blender/depsgraph/intern/builder/deg_builder_relations.h b/source/blender/depsgraph/intern/builder/deg_builder_relations.h index d666e59d368..7a7b3f7b86c 100644 --- a/source/blender/depsgraph/intern/builder/deg_builder_relations.h +++ b/source/blender/depsgraph/intern/builder/deg_builder_relations.h @@ -249,6 +249,8 @@ struct DepsgraphRelationBuilder template <typename KeyType> OperationDepsNode *find_operation_node(const KeyType &key); + Depsgraph *getGraph(); + protected: RootDepsNode *find_node(const RootKey &key) const; TimeSourceDepsNode *find_node(const TimeSourceKey &key) const; diff --git a/source/blender/depsgraph/intern/depsgraph_build.cc b/source/blender/depsgraph/intern/depsgraph_build.cc index 82be671236f..f538f51b272 100644 --- a/source/blender/depsgraph/intern/depsgraph_build.cc +++ b/source/blender/depsgraph/intern/depsgraph_build.cc @@ -166,6 +166,13 @@ void DEG_add_bone_relation(DepsNodeHandle *handle, description); } +struct Depsgraph *DEG_get_graph_from_handle(struct DepsNodeHandle *handle) +{ + DEG::DepsNodeHandle *deg_handle = get_handle(handle); + DEG::DepsgraphRelationBuilder *relation_builder = deg_handle->builder; + return reinterpret_cast<Depsgraph *>(relation_builder->getGraph()); +} + void DEG_add_special_eval_flag(Depsgraph *graph, ID *id, short flag) { DEG::Depsgraph *deg_graph = reinterpret_cast<DEG::Depsgraph *>(graph); diff --git a/source/blender/editors/animation/keyframes_edit.c b/source/blender/editors/animation/keyframes_edit.c index f4922fec385..5d1366e03f8 100644 --- a/source/blender/editors/animation/keyframes_edit.c +++ b/source/blender/editors/animation/keyframes_edit.c @@ -452,7 +452,7 @@ void ANIM_editkeyframes_refresh(bAnimContext *ac) ok |= KEYFRAME_OK_H2; \ } \ } (void)0 - + /* ------------------------ */ static short ok_bezier_frame(KeyframeEditData *ked, BezTriple *bezt) diff --git a/source/blender/editors/include/UI_icons.h b/source/blender/editors/include/UI_icons.h index 8420591aa3e..0c83038b7a3 100644 --- a/source/blender/editors/include/UI_icons.h +++ b/source/blender/editors/include/UI_icons.h @@ -313,9 +313,9 @@ DEF_ICON(OUTLINER_OB_ARMATURE) DEF_ICON(OUTLINER_OB_FONT) DEF_ICON(OUTLINER_OB_SURFACE) DEF_ICON(OUTLINER_OB_SPEAKER) +DEF_ICON(OUTLINER_OB_FORCE_FIELD) +DEF_ICON(OUTLINER_OB_GROUP_INSTANCE) #ifndef DEF_ICON_BLANK_SKIP - DEF_ICON(BLANK120) - DEF_ICON(BLANK121) DEF_ICON(BLANK122) DEF_ICON(BLANK123) DEF_ICON(BLANK124) diff --git a/source/blender/editors/io/io_alembic.c b/source/blender/editors/io/io_alembic.c index ba8792d12ff..fb20d9f3caa 100644 --- a/source/blender/editors/io/io_alembic.c +++ b/source/blender/editors/io/io_alembic.c @@ -383,8 +383,8 @@ void WM_OT_alembic_export(wmOperatorType *ot) "Enable this to run the import in the background, disable to block Blender while importing"); /* This dummy prop is used to check whether we need to init the start and - * end frame values to that of the scene's, otherwise they are reset at - * every change, draw update. */ + * end frame values to that of the scene's, otherwise they are reset at + * every change, draw update. */ RNA_def_boolean(ot->srna, "init_scene_frame_range", false, "", ""); } diff --git a/source/blender/editors/object/object_relations.c b/source/blender/editors/object/object_relations.c index ae3a3866c24..333bcf00a8c 100644 --- a/source/blender/editors/object/object_relations.c +++ b/source/blender/editors/object/object_relations.c @@ -2055,24 +2055,6 @@ void ED_object_single_users(Main *bmain, Scene *scene, const bool full, const bo /******************************* Make Local ***********************************/ -/* helper for below, ma was checked to be not NULL */ -static void make_local_makelocalmaterial(Material *ma) -{ - AnimData *adt; - int b; - - id_make_local(G.main, &ma->id, false, false); - - for (b = 0; b < MAX_MTEX; b++) - if (ma->mtex[b] && ma->mtex[b]->tex) - id_make_local(G.main, &ma->mtex[b]->tex->id, false, false); - - adt = BKE_animdata_from_id(&ma->id); - if (adt) BKE_animdata_make_local(adt); - - /* nodetree? XXX */ -} - enum { MAKE_LOCAL_SELECT_OB = 1, MAKE_LOCAL_SELECT_OBDATA = 2, @@ -2159,125 +2141,145 @@ static bool make_local_all__instance_indirect_unused(Main *bmain, Scene *scene, return changed; } -static int make_local_exec(bContext *C, wmOperator *op) +static void make_local_animdata_tag_strips(ListBase *strips) { - Main *bmain = CTX_data_main(C); - Scene *scene = CTX_data_scene(C); - SceneLayer *sl = CTX_data_scene_layer(C); - SceneCollection *sc = CTX_data_scene_collection(C); - AnimData *adt; - ParticleSystem *psys; - Material *ma, ***matarar; - Lamp *la; - ID *id; - const int mode = RNA_enum_get(op->ptr, "type"); - int a, b; - - if (mode == MAKE_LOCAL_ALL) { - /* de-select so the user can differentiate newly instanced from existing objects */ - BKE_scene_base_deselect_all(scene); + NlaStrip *strip; - if (make_local_all__instance_indirect_unused(bmain, scene, sl, sc)) { - BKE_report(op->reports, RPT_INFO, - "Orphan library objects added to the current scene to avoid loss"); + for (strip = strips->first; strip; strip = strip->next) { + if (strip->act) { + strip->act->id.tag &= ~LIB_TAG_PRE_EXISTING; + } + if (strip->remap && strip->remap->target) { + strip->remap->target->id.tag &= ~LIB_TAG_PRE_EXISTING; } - BKE_library_make_local(bmain, NULL, NULL, false, false); /* NULL is all libs */ - WM_event_add_notifier(C, NC_WINDOW, NULL); - return OPERATOR_FINISHED; + make_local_animdata_tag_strips(&strip->strips); } +} - tag_localizable_objects(C, mode); - - CTX_DATA_BEGIN (C, Object *, ob, selected_objects) - { - if ((ob->id.tag & LIB_TAG_DOIT) == 0) { - continue; +/* Tag all actions used by given animdata to be made local. */ +static void make_local_animdata_tag(AnimData *adt) +{ + if (adt) { + /* Actions - Active and Temp */ + if (adt->action) { + adt->action->id.tag &= ~LIB_TAG_PRE_EXISTING; + } + if (adt->tmpact) { + adt->tmpact->id.tag &= ~LIB_TAG_PRE_EXISTING; + } + /* Remaps */ + if (adt->remap && adt->remap->target) { + adt->remap->target->id.tag &= ~LIB_TAG_PRE_EXISTING; } - if (ob->id.lib) - id_make_local(bmain, &ob->id, false, false); - } - CTX_DATA_END; + /* Drivers */ + /* TODO: need to handle the ID-targets too? */ - /* maybe object pointers */ - CTX_DATA_BEGIN (C, Object *, ob, selected_objects) - { - if (ob->id.lib == NULL) { - ID_NEW_REMAP(ob->parent); + /* NLA Data */ + for (NlaTrack *nlt = adt->nla_tracks.first; nlt; nlt = nlt->next) { + make_local_animdata_tag_strips(&nlt->strips); } } - CTX_DATA_END; - - CTX_DATA_BEGIN (C, Object *, ob, selected_objects) - { - if ((ob->id.tag & LIB_TAG_DOIT) == 0) { - continue; - } +} - id = ob->data; +static void make_local_material_tag(Material *ma) +{ + if (ma) { + ma->id.tag &= ~LIB_TAG_PRE_EXISTING; + make_local_animdata_tag(BKE_animdata_from_id(&ma->id)); - if (id && (ELEM(mode, MAKE_LOCAL_SELECT_OBDATA, MAKE_LOCAL_SELECT_OBDATA_MATERIAL))) { - id_make_local(bmain, id, false, false); - adt = BKE_animdata_from_id(id); - if (adt) BKE_animdata_make_local(adt); + /* About nodetrees: root one is made local together with material, others we keep linked for now... */ - /* tag indirect data direct */ - matarar = give_matarar(ob); - if (matarar) { - for (a = 0; a < ob->totcol; a++) { - ma = (*matarar)[a]; - if (ma) - id_lib_extern(&ma->id); - } + for (int a = 0; a < MAX_MTEX; a++) { + if (ma->mtex[a] && ma->mtex[a]->tex) { + ma->mtex[a]->tex->id.tag &= ~LIB_TAG_PRE_EXISTING; } } + } +} - for (psys = ob->particlesystem.first; psys; psys = psys->next) - id_make_local(bmain, &psys->part->id, false, false); +static int make_local_exec(bContext *C, wmOperator *op) +{ + Main *bmain = CTX_data_main(C); + Scene *scene = CTX_data_scene(C); + ParticleSystem *psys; + Material *ma, ***matarar; + Lamp *la; + const int mode = RNA_enum_get(op->ptr, "type"); + int a; + + /* Note: we (ab)use LIB_TAG_PRE_EXISTING to cherry pick which ID to make local... */ + if (mode == MAKE_LOCAL_ALL) { + SceneLayer *sl = CTX_data_scene_layer(C); + SceneCollection *sc = CTX_data_scene_collection(C); + + BKE_main_id_tag_all(bmain, LIB_TAG_PRE_EXISTING, false); + + /* de-select so the user can differentiate newly instanced from existing objects */ + BKE_scene_base_deselect_all(scene); - adt = BKE_animdata_from_id(&ob->id); - if (adt) BKE_animdata_make_local(adt); + if (make_local_all__instance_indirect_unused(bmain, scene, sl, sc)) { + BKE_report(op->reports, RPT_INFO, "Orphan library objects added to the current scene to avoid loss"); + } } - CTX_DATA_END; + else { + BKE_main_id_tag_all(bmain, LIB_TAG_PRE_EXISTING, true); + tag_localizable_objects(C, mode); - if (mode == MAKE_LOCAL_SELECT_OBDATA_MATERIAL) { CTX_DATA_BEGIN (C, Object *, ob, selected_objects) { if ((ob->id.tag & LIB_TAG_DOIT) == 0) { continue; } - if (ob->type == OB_LAMP) { - la = ob->data; - - for (b = 0; b < MAX_MTEX; b++) - if (la->mtex[b] && la->mtex[b]->tex) - id_make_local(bmain, &la->mtex[b]->tex->id, false, false); + ob->id.tag &= ~LIB_TAG_PRE_EXISTING; + make_local_animdata_tag(BKE_animdata_from_id(&ob->id)); + for (psys = ob->particlesystem.first; psys; psys = psys->next) { + psys->part->id.tag &= ~LIB_TAG_PRE_EXISTING; } - else { + + if (mode == MAKE_LOCAL_SELECT_OBDATA_MATERIAL) { for (a = 0; a < ob->totcol; a++) { ma = ob->mat[a]; - if (ma) - make_local_makelocalmaterial(ma); + if (ma) { + make_local_material_tag(ma); + } } matarar = (Material ***)give_matarar(ob); if (matarar) { for (a = 0; a < ob->totcol; a++) { ma = (*matarar)[a]; - if (ma) - make_local_makelocalmaterial(ma); + if (ma) { + make_local_material_tag(ma); + } + } + } + + if (ob->type == OB_LAMP) { + BLI_assert(ob->data != NULL); + la = ob->data; + for (a = 0; a < MAX_MTEX; a++) { + if (la->mtex[a] && la->mtex[a]->tex) { + la->id.tag &= ~LIB_TAG_PRE_EXISTING; + } } } } + + if (ELEM(mode, MAKE_LOCAL_SELECT_OBDATA, MAKE_LOCAL_SELECT_OBDATA_MATERIAL) && ob->data != NULL) { + ID *ob_data = ob->data; + ob_data->tag &= ~LIB_TAG_PRE_EXISTING; + make_local_animdata_tag(BKE_animdata_from_id(ob_data)); + } } CTX_DATA_END; } - BKE_main_id_clear_newpoins(bmain); - WM_event_add_notifier(C, NC_WINDOW, NULL); + BKE_library_make_local(bmain, NULL, NULL, true, false); /* NULL is all libs */ + WM_event_add_notifier(C, NC_WINDOW, NULL); return OPERATOR_FINISHED; } diff --git a/source/blender/editors/physics/dynamicpaint_ops.c b/source/blender/editors/physics/dynamicpaint_ops.c index a5e78e116c6..cd07b57577b 100644 --- a/source/blender/editors/physics/dynamicpaint_ops.c +++ b/source/blender/editors/physics/dynamicpaint_ops.c @@ -281,10 +281,10 @@ void DPAINT_OT_output_toggle(wmOperatorType *ot) /***************************** Image Sequence Baking ******************************/ typedef struct DynamicPaintBakeJob { - /* from wmJob */ - void *owner; - short *stop, *do_update; - float *progress; + /* from wmJob */ + void *owner; + short *stop, *do_update; + float *progress; struct Main *bmain; Scene *scene; @@ -300,13 +300,13 @@ typedef struct DynamicPaintBakeJob { static void dpaint_bake_free(void *customdata) { - DynamicPaintBakeJob *job = customdata; - MEM_freeN(job); + DynamicPaintBakeJob *job = customdata; + MEM_freeN(job); } static void dpaint_bake_endjob(void *customdata) { - DynamicPaintBakeJob *job = customdata; + DynamicPaintBakeJob *job = customdata; DynamicPaintCanvasSettings *canvas = job->canvas; canvas->flags &= ~MOD_DPAINT_BAKING; @@ -314,7 +314,7 @@ static void dpaint_bake_endjob(void *customdata) dynamicPaint_freeSurfaceData(job->surface); G.is_rendering = false; - BKE_spacedata_draw_locks(false); + BKE_spacedata_draw_locks(false); WM_set_locked_interface(G.main->wm.first, false); @@ -424,26 +424,26 @@ static void dynamicPaint_bakeImageSequence(DynamicPaintBakeJob *job) static void dpaint_bake_startjob(void *customdata, short *stop, short *do_update, float *progress) { - DynamicPaintBakeJob *job = customdata; + DynamicPaintBakeJob *job = customdata; - job->stop = stop; - job->do_update = do_update; - job->progress = progress; + job->stop = stop; + job->do_update = do_update; + job->progress = progress; job->start = PIL_check_seconds_timer(); job->success = 1; - G.is_break = false; /* reset BKE_blender_test_break*/ + G.is_break = false; /* reset BKE_blender_test_break*/ /* XXX annoying hack: needed to prevent data corruption when changing * scene frame in separate threads - */ - G.is_rendering = true; - BKE_spacedata_draw_locks(true); + */ + G.is_rendering = true; + BKE_spacedata_draw_locks(true); dynamicPaint_bakeImageSequence(job); - *do_update = true; - *stop = 0; + *do_update = true; + *stop = 0; } /* diff --git a/source/blender/editors/physics/physics_pointcache.c b/source/blender/editors/physics/physics_pointcache.c index 5a2bb1e74c7..d4f949880ee 100644 --- a/source/blender/editors/physics/physics_pointcache.c +++ b/source/blender/editors/physics/physics_pointcache.c @@ -100,45 +100,45 @@ static int ptcache_job_break(void *customdata) static void ptcache_job_update(void *customdata, float progress, int *cancel) { - PointCacheJob *job = customdata; + PointCacheJob *job = customdata; - if (ptcache_job_break(job)) { - *cancel = 1; - } + if (ptcache_job_break(job)) { + *cancel = 1; + } - *(job->do_update) = true; - *(job->progress) = progress; + *(job->do_update) = true; + *(job->progress) = progress; } static void ptcache_job_startjob(void *customdata, short *stop, short *do_update, float *progress) { - PointCacheJob *job = customdata; + PointCacheJob *job = customdata; - job->stop = stop; - job->do_update = do_update; - job->progress = progress; + job->stop = stop; + job->do_update = do_update; + job->progress = progress; - G.is_break = false; + G.is_break = false; - /* XXX annoying hack: needed to prevent data corruption when changing - * scene frame in separate threads - */ - G.is_rendering = true; - BKE_spacedata_draw_locks(true); + /* XXX annoying hack: needed to prevent data corruption when changing + * scene frame in separate threads + */ + G.is_rendering = true; + BKE_spacedata_draw_locks(true); BKE_ptcache_bake(job->baker); - *do_update = true; - *stop = 0; + *do_update = true; + *stop = 0; } static void ptcache_job_endjob(void *customdata) { - PointCacheJob *job = customdata; + PointCacheJob *job = customdata; Scene *scene = job->baker->scene; - G.is_rendering = false; - BKE_spacedata_draw_locks(false); + G.is_rendering = false; + BKE_spacedata_draw_locks(false); WM_set_locked_interface(G.main->wm.first, false); diff --git a/source/blender/editors/space_nla/nla_buttons.c b/source/blender/editors/space_nla/nla_buttons.c index 5355b8012db..c774b99629c 100644 --- a/source/blender/editors/space_nla/nla_buttons.c +++ b/source/blender/editors/space_nla/nla_buttons.c @@ -552,7 +552,7 @@ void nla_buttons_register(ARegionType *art) pt = MEM_callocN(sizeof(PanelType), "spacetype nla panel modifiers"); strcpy(pt->idname, "NLA_PT_modifiers"); strcpy(pt->label, N_("Modifiers")); - strcpy(pt->category, "Modifiers"); + strcpy(pt->category, "Modifiers"); strcpy(pt->translation_context, BLT_I18NCONTEXT_DEFAULT_BPYRNA); pt->draw = nla_panel_modifiers; pt->poll = nla_strip_eval_panel_poll; diff --git a/source/blender/editors/space_sequencer/sequencer_select.c b/source/blender/editors/space_sequencer/sequencer_select.c index 48c49f36471..d88ed36e392 100644 --- a/source/blender/editors/space_sequencer/sequencer_select.c +++ b/source/blender/editors/space_sequencer/sequencer_select.c @@ -683,7 +683,7 @@ static int sequencer_select_less_exec(bContext *C, wmOperator *UNUSED(op)) if (!select_more_less_seq__internal(scene, false, false)) return OPERATOR_CANCELLED; - + WM_event_add_notifier(C, NC_SCENE | ND_SEQUENCER | NA_SELECTED, scene); return OPERATOR_FINISHED; diff --git a/source/blender/editors/transform/transform_generics.c b/source/blender/editors/transform/transform_generics.c index e0f295b3696..3d80ac11ec9 100644 --- a/source/blender/editors/transform/transform_generics.c +++ b/source/blender/editors/transform/transform_generics.c @@ -975,7 +975,7 @@ static void recalcData_sequencer(TransInfo *t) /* force recalculation of triangles during transformation */ static void recalcData_gpencil_strokes(TransInfo *t) - { +{ TransData *td = t->data; for (int i = 0; i < t->total; i++, td++) { bGPDstroke *gps = td->extra; diff --git a/source/blender/editors/uvedit/uvedit_parametrizer.c b/source/blender/editors/uvedit/uvedit_parametrizer.c index bdfff123aa4..8c76d03035a 100644 --- a/source/blender/editors/uvedit/uvedit_parametrizer.c +++ b/source/blender/editors/uvedit/uvedit_parametrizer.c @@ -2859,7 +2859,7 @@ static PBool p_chart_symmetry_pins(PChart *chart, PEdge *outer, PVert **pin1, PV PEdge *cure = NULL, *firste1 = NULL, *firste2 = NULL, *nextbe; float maxlen = 0.0f, curlen = 0.0f, totlen = 0.0f, firstlen = 0.0f; float len1, len2; - + /* find longest series of verts split in the chart itself, these are * marked during construction */ be = outer; diff --git a/source/blender/makesrna/intern/rna_access.c b/source/blender/makesrna/intern/rna_access.c index 8d7ad46630c..39e115c1891 100644 --- a/source/blender/makesrna/intern/rna_access.c +++ b/source/blender/makesrna/intern/rna_access.c @@ -6450,7 +6450,7 @@ static int rna_function_parameter_parse(PointerRNA *ptr, PropertyRNA *prop, Prop tid, fid, pid, RNA_struct_identifier(ptype), RNA_struct_identifier(srna)); return -1; } - + *((void **)dest) = *((void **)src); break; diff --git a/source/blender/makesrna/intern/rna_particle.c b/source/blender/makesrna/intern/rna_particle.c index 53c40996b51..29f33907d39 100644 --- a/source/blender/makesrna/intern/rna_particle.c +++ b/source/blender/makesrna/intern/rna_particle.c @@ -2097,7 +2097,7 @@ static void rna_def_particle_settings(BlenderRNA *brna) RNA_def_property_ui_text(prop, "Even Distribution", "Use even distribution from faces based on face areas or edge lengths"); RNA_def_property_update(prop, 0, "rna_Particle_reset"); - + prop = RNA_def_property(srna, "use_die_on_collision", PROP_BOOLEAN, PROP_NONE); RNA_def_property_boolean_sdna(prop, NULL, "flag", PART_DIE_ON_COL); RNA_def_property_clear_flag(prop, PROP_ANIMATABLE); diff --git a/source/blender/makesrna/intern/rna_userdef.c b/source/blender/makesrna/intern/rna_userdef.c index ae77f64497e..1d01e85205c 100644 --- a/source/blender/makesrna/intern/rna_userdef.c +++ b/source/blender/makesrna/intern/rna_userdef.c @@ -3318,7 +3318,7 @@ static void rna_def_userdef_view(BlenderRNA *brna) prop = RNA_def_property(srna, "ui_scale", PROP_FLOAT, PROP_FACTOR); RNA_def_property_ui_text(prop, "UI Scale", "Changes the size of the fonts and buttons in the interface"); RNA_def_property_range(prop, 0.25f, 4.0f); - RNA_def_property_ui_range(prop, 0.5f, 2.0f, 1, 1); + RNA_def_property_ui_range(prop, 0.5f, 2.0f, 1, 2); RNA_def_property_float_default(prop, 1.0f); RNA_def_property_update(prop, 0, "rna_userdef_dpi_update"); diff --git a/source/blender/modifiers/intern/MOD_array.c b/source/blender/modifiers/intern/MOD_array.c index 09597b5b42f..57f90fb4b51 100644 --- a/source/blender/modifiers/intern/MOD_array.c +++ b/source/blender/modifiers/intern/MOD_array.c @@ -104,7 +104,7 @@ static void foreachObjectLink( static void updateDepsgraph(ModifierData *md, struct Main *UNUSED(bmain), - struct Scene *scene, + struct Scene *UNUSED(scene), Object *UNUSED(ob), struct DepsNodeHandle *node) { @@ -116,8 +116,9 @@ static void updateDepsgraph(ModifierData *md, DEG_add_object_relation(node, amd->end_cap, DEG_OB_COMP_TRANSFORM, "Array Modifier End Cap"); } if (amd->curve_ob) { + struct Depsgraph *depsgraph = DEG_get_graph_from_handle(node); DEG_add_object_relation(node, amd->curve_ob, DEG_OB_COMP_GEOMETRY, "Array Modifier Curve"); - DEG_add_special_eval_flag(scene->depsgraph, &amd->curve_ob->id, DAG_EVAL_NEED_CURVE_PATH); + DEG_add_special_eval_flag(depsgraph, &amd->curve_ob->id, DAG_EVAL_NEED_CURVE_PATH); } if (amd->offset_ob != NULL) { DEG_add_object_relation(node, amd->offset_ob, DEG_OB_COMP_TRANSFORM, "Array Modifier Offset"); diff --git a/source/blender/modifiers/intern/MOD_curve.c b/source/blender/modifiers/intern/MOD_curve.c index d9acce0a780..3200caedc32 100644 --- a/source/blender/modifiers/intern/MOD_curve.c +++ b/source/blender/modifiers/intern/MOD_curve.c @@ -93,7 +93,7 @@ static void foreachObjectLink( static void updateDepsgraph(ModifierData *md, struct Main *UNUSED(bmain), - struct Scene *scene, + struct Scene *UNUSED(scene), Object *object, struct DepsNodeHandle *node) { @@ -105,8 +105,9 @@ static void updateDepsgraph(ModifierData *md, /* TODO(sergey): Currently path is evaluated as a part of modifier stack, * might be changed in the future. */ + struct Depsgraph *depsgraph = DEG_get_graph_from_handle(node); DEG_add_object_relation(node, cmd->object, DEG_OB_COMP_GEOMETRY, "Curve Modifier"); - DEG_add_special_eval_flag(scene->depsgraph, &cmd->object->id, DAG_EVAL_NEED_CURVE_PATH); + DEG_add_special_eval_flag(depsgraph, &cmd->object->id, DAG_EVAL_NEED_CURVE_PATH); } DEG_add_object_relation(node, object, DEG_OB_COMP_TRANSFORM, "Curve Modifier"); diff --git a/source/blender/nodes/shader/nodes/node_shader_normal_map.c b/source/blender/nodes/shader/nodes/node_shader_normal_map.c index e0bf34f42e4..36d7522e3e6 100644 --- a/source/blender/nodes/shader/nodes/node_shader_normal_map.c +++ b/source/blender/nodes/shader/nodes/node_shader_normal_map.c @@ -46,8 +46,10 @@ static void node_shader_init_normal_map(bNodeTree *UNUSED(ntree), bNode *node) node->storage = attr; } -static void node_shader_exec_normal_map(void *data, int UNUSED(thread), bNode *node, bNodeExecData *UNUSED(execdata), bNodeStack **in, bNodeStack **out) - { +static void node_shader_exec_normal_map( + void *data, int UNUSED(thread), bNode *node, bNodeExecData *UNUSED(execdata), + bNodeStack **in, bNodeStack **out) +{ if (data) { ShadeInput *shi = ((ShaderCallData *)data)->shi; diff --git a/source/blender/python/intern/bpy_app_sdl.c b/source/blender/python/intern/bpy_app_sdl.c index 2f4d8e6c325..76dab775953 100644 --- a/source/blender/python/intern/bpy_app_sdl.c +++ b/source/blender/python/intern/bpy_app_sdl.c @@ -56,7 +56,7 @@ static PyStructSequence_Field app_sdl_info_fields[] = { {(char *)"available", (char *)("Boolean, True when SDL is available. This is False when " "either *supported* is False, or *dynload* is True and " "Blender cannot find the correct library.")}, - {NULL} + {NULL} }; static PyStructSequence_Desc app_sdl_info_desc = { diff --git a/source/blender/render/intern/source/sunsky.c b/source/blender/render/intern/source/sunsky.c index d4e53eb7305..f0cf29e98ca 100644 --- a/source/blender/render/intern/source/sunsky.c +++ b/source/blender/render/intern/source/sunsky.c @@ -398,7 +398,7 @@ void InitAtmosphere(struct SunSky *sunSky, float sun_intens, float mief, float r vLambda2[0] = fLambda2[0]; vLambda2[1] = fLambda2[1]; vLambda2[2] = fLambda2[2]; - + vLambda4[0] = fLambda4[0]; vLambda4[1] = fLambda4[1]; vLambda4[2] = fLambda4[2]; diff --git a/source/blender/windowmanager/intern/wm_files.c b/source/blender/windowmanager/intern/wm_files.c index 8d442a2ed25..4e8a0cf8d05 100644 --- a/source/blender/windowmanager/intern/wm_files.c +++ b/source/blender/windowmanager/intern/wm_files.c @@ -1100,7 +1100,7 @@ static int wm_file_write(bContext *C, const char *filepath, int fileflags, Repor BKE_reportf(reports, RPT_ERROR, "Cannot save blend file, path '%s' is not writable", filepath); return ret; } - + /* note: used to replace the file extension (to ensure '.blend'), * no need to now because the operator ensures, * its handy for scripts to save to a predefined name without blender editing it */ diff --git a/tests/gtests/blenlib/BLI_kdopbvh_test.cc b/tests/gtests/blenlib/BLI_kdopbvh_test.cc new file mode 100644 index 00000000000..bc8cf58d69d --- /dev/null +++ b/tests/gtests/blenlib/BLI_kdopbvh_test.cc @@ -0,0 +1,92 @@ +/* Apache License, Version 2.0 */ + +#include "testing/testing.h" + +/* TODO: ray intersection, overlap ... etc.*/ + +extern "C" { +#include "BLI_compiler_attrs.h" +#include "BLI_kdopbvh.h" +#include "BLI_rand.h" +#include "BLI_math_vector.h" +#include "MEM_guardedalloc.h" +} + +/* -------------------------------------------------------------------- */ +/* Helper Functions */ + +static void rng_v3_round( + float *coords, int coords_len, + struct RNG *rng, int round, float scale) +{ + for (int i = 0; i < coords_len; i++) { + float f = BLI_rng_get_float(rng) * 2.0f - 1.0f; + coords[i] = ((float)((int)(f * round)) / (float)round) * scale; + } +} + +/* -------------------------------------------------------------------- */ +/* Tests */ + +TEST(kdopbvh, Empty) +{ + BVHTree *tree = BLI_bvhtree_new(0, 0.0, 8, 8); + BLI_bvhtree_balance(tree); + EXPECT_EQ(0, BLI_bvhtree_get_size(tree)); + BLI_bvhtree_free(tree); +} + +TEST(kdopbvh, Single) +{ + BVHTree *tree = BLI_bvhtree_new(1, 0.0, 8, 8); + { + float co[3] = {0}; + BLI_bvhtree_insert(tree, 0, co, 1); + } + + EXPECT_EQ(BLI_bvhtree_get_size(tree), 1); + + BLI_bvhtree_balance(tree); + BLI_bvhtree_free(tree); +} + +/** + * Note that a small epsilon is added to the BVH nodes bounds, even if we pass in zero. + * Use rounding to ensure very close nodes don't cause the wrong node to be found as nearest. + */ +static void find_nearest_points_test(int points_len, float scale, int round, int random_seed) +{ + struct RNG *rng = BLI_rng_new(random_seed); + BVHTree *tree = BLI_bvhtree_new(points_len, 0.0, 8, 8); + float (*points)[3] = (float (*)[3])MEM_mallocN(sizeof(float[3]) * points_len, __func__); + for (int i = 0; i < points_len; i++) { + rng_v3_round(points[i], 3, rng, round, scale); + BLI_bvhtree_insert(tree, i, points[i], 1); + } + BLI_bvhtree_balance(tree); + /* first find each point */ + for (int i = 0; i < points_len; i++) { + const int j = BLI_bvhtree_find_nearest(tree, points[i], NULL, NULL, NULL); + if (j != i) { +#if 0 + const float dist = len_v3v3(points[i], points[j]); + if (dist > (1.0f / (float)round)) { + printf("%.15f (%d %d)\n", dist, i, j); + print_v3_id(points[i]); + print_v3_id(points[j]); + fflush(stdout); + } +#endif + EXPECT_GE(j, 0); + EXPECT_LT(j, points_len); + EXPECT_EQ_ARRAY(points[i], points[j], 3); + } + } + BLI_bvhtree_free(tree); + BLI_rng_free(rng); + MEM_freeN(points); +} + +TEST(kdopbvh, FindNearest_1) { find_nearest_points_test(1, 1.0, 1000, 1234); } +TEST(kdopbvh, FindNearest_2) { find_nearest_points_test(2, 1.0, 1000, 123); } +TEST(kdopbvh, FindNearest_500) { find_nearest_points_test(500, 1.0, 1000, 12); } diff --git a/tests/gtests/blenlib/CMakeLists.txt b/tests/gtests/blenlib/CMakeLists.txt index a190d9cd8c5..8b013e7a7a6 100644 --- a/tests/gtests/blenlib/CMakeLists.txt +++ b/tests/gtests/blenlib/CMakeLists.txt @@ -37,6 +37,7 @@ set(CMAKE_EXE_LINKER_FLAGS_DEBUG "${CMAKE_EXE_LINKER_FLAGS_DEBUG} ${PLATFORM_LIN BLENDER_TEST(BLI_array_store "bf_blenlib") BLENDER_TEST(BLI_array_utils "bf_blenlib") +BLENDER_TEST(BLI_kdopbvh "bf_blenlib;bf_intern_eigen") BLENDER_TEST(BLI_stack "bf_blenlib") BLENDER_TEST(BLI_math_color "bf_blenlib") BLENDER_TEST(BLI_math_geom "bf_blenlib;bf_intern_eigen") diff --git a/tests/python/bl_alembic_import_test.py b/tests/python/bl_alembic_import_test.py index 1322a883b72..e9dcbd80a3a 100644 --- a/tests/python/bl_alembic_import_test.py +++ b/tests/python/bl_alembic_import_test.py @@ -207,7 +207,7 @@ def main(): import argparse if '--' in sys.argv: - argv = [sys.argv[0]] + sys.argv[sys.argv.index('--')+1:] + argv = [sys.argv[0]] + sys.argv[sys.argv.index('--') + 1:] else: argv = sys.argv diff --git a/tests/python/bl_keymap_completeness.py b/tests/python/bl_keymap_completeness.py index 00322907f69..652ed449a3c 100644 --- a/tests/python/bl_keymap_completeness.py +++ b/tests/python/bl_keymap_completeness.py @@ -80,5 +80,6 @@ def main(): import sys sys.exit(1) + if __name__ == "__main__": main() diff --git a/tests/python/pep8.py b/tests/python/pep8.py index 0e6250f534b..dde4250f6aa 100644 --- a/tests/python/pep8.py +++ b/tests/python/pep8.py @@ -178,5 +178,6 @@ def main(): "--max-line-length=1000" " '%s'" % f) + if __name__ == "__main__": main() diff --git a/tests/python/rna_info_dump.py b/tests/python/rna_info_dump.py index c26d94a1246..da228e52652 100644 --- a/tests/python/rna_info_dump.py +++ b/tests/python/rna_info_dump.py @@ -127,5 +127,6 @@ def api_dump(use_properties=True, use_functions=True): print("END") + if __name__ == "__main__": api_dump() |
