Merge branch 'master' into soc-2017-sculpting_improvements
authorwitt <sebastian.witt@rwth-aachen.de>
Mon, 12 Jun 2017 08:12:49 +0000 (10:12 +0200)
committerwitt <sebastian.witt@rwth-aachen.de>
Mon, 12 Jun 2017 08:12:49 +0000 (10:12 +0200)
Keep the branch updated to make merging it easier later.

441 files changed:
CMakeLists.txt
build_files/build_environment/install_deps.sh
build_files/buildbot/master_unpack.py
build_files/cmake/Modules/FindOpenCOLLADA.cmake
build_files/cmake/config/blender_full.cmake
build_files/cmake/config/blender_release.cmake
build_files/cmake/macros.cmake
build_files/cmake/platform/platform_win32.cmake
build_files/cmake/platform/platform_win32_mingw.cmake [deleted file]
build_files/cmake/platform/platform_win32_msvc.cmake
build_files/cmake/project_info.py
doc/python_api/examples/bpy.types.Operator.5.py
doc/python_api/rst/info_overview.rst
extern/carve/include/carve/win32.h
extern/carve/patches/mingw.patch [deleted file]
extern/carve/patches/mingw_w64.patch [deleted file]
extern/carve/patches/series
extern/carve/patches/win32.patch
intern/audaspace/CMakeLists.txt
intern/cycles/blender/addon/engine.py
intern/cycles/blender/addon/properties.py
intern/cycles/blender/addon/ui.py
intern/cycles/blender/blender_mesh.cpp
intern/cycles/blender/blender_session.cpp
intern/cycles/blender/blender_sync.cpp
intern/cycles/blender/blender_sync.h
intern/cycles/blender/blender_util.h
intern/cycles/device/device.cpp
intern/cycles/device/device.h
intern/cycles/device/device_cpu.cpp
intern/cycles/device/device_cuda.cpp
intern/cycles/device/device_denoising.cpp
intern/cycles/device/device_denoising.h
intern/cycles/device/device_opencl.cpp
intern/cycles/device/device_split_kernel.cpp
intern/cycles/device/device_split_kernel.h
intern/cycles/device/opencl/opencl.h
intern/cycles/device/opencl/opencl_base.cpp
intern/cycles/device/opencl/opencl_split.cpp
intern/cycles/device/opencl/opencl_util.cpp
intern/cycles/kernel/CMakeLists.txt
intern/cycles/kernel/closure/bsdf_microfacet_multi.h
intern/cycles/kernel/closure/bsdf_microfacet_multi_impl.h
intern/cycles/kernel/filter/filter_features.h
intern/cycles/kernel/filter/filter_features_sse.h
intern/cycles/kernel/filter/filter_nlm_cpu.h
intern/cycles/kernel/filter/filter_nlm_gpu.h
intern/cycles/kernel/filter/filter_prefilter.h
intern/cycles/kernel/filter/filter_reconstruction.h
intern/cycles/kernel/filter/filter_transform.h
intern/cycles/kernel/filter/filter_transform_gpu.h
intern/cycles/kernel/filter/filter_transform_sse.h
intern/cycles/kernel/kernel_accumulate.h
intern/cycles/kernel/kernel_compat_cpu.h
intern/cycles/kernel/kernel_compat_cuda.h
intern/cycles/kernel/kernel_compat_opencl.h
intern/cycles/kernel/kernel_passes.h
intern/cycles/kernel/kernel_path_state.h
intern/cycles/kernel/kernel_queues.h
intern/cycles/kernel/kernel_types.h
intern/cycles/kernel/kernels/cpu/filter_cpu.h
intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
intern/cycles/kernel/kernels/cpu/kernel_cpu.h
intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
intern/cycles/kernel/kernels/cuda/filter.cu
intern/cycles/kernel/kernels/cuda/kernel_split.cu
intern/cycles/kernel/kernels/opencl/filter.cl
intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl [new file with mode: 0644]
intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl
intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl
intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
intern/cycles/kernel/kernels/opencl/kernel_split.cl
intern/cycles/kernel/kernels/opencl/kernel_split_function.h [new file with mode: 0644]
intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl
intern/cycles/kernel/osl/osl_bssrdf.cpp
intern/cycles/kernel/shaders/node_principled_bsdf.osl
intern/cycles/kernel/split/kernel_branched.h
intern/cycles/kernel/split/kernel_buffer_update.h
intern/cycles/kernel/split/kernel_do_volume.h
intern/cycles/kernel/split/kernel_enqueue_inactive.h [new file with mode: 0644]
intern/cycles/kernel/split/kernel_next_iteration_setup.h
intern/cycles/kernel/split/kernel_scene_intersect.h
intern/cycles/kernel/split/kernel_shadow_blocked_dl.h
intern/cycles/kernel/split/kernel_split_common.h
intern/cycles/kernel/split/kernel_split_data_types.h
intern/cycles/kernel/split/kernel_subsurface_scatter.h
intern/cycles/kernel/svm/svm_closure.h
intern/cycles/render/image.cpp
intern/cycles/render/nodes.cpp
intern/cycles/render/nodes.h
intern/cycles/render/session.cpp
intern/cycles/render/session.h
intern/cycles/render/shader.cpp
intern/cycles/render/shader.h
intern/cycles/test/util_string_test.cpp
intern/cycles/util/util_atomic.h
intern/cycles/util/util_guarded_allocator.h
intern/cycles/util/util_math_float3.h
intern/cycles/util/util_math_matrix.h
intern/cycles/util/util_progress.h
intern/cycles/util/util_string.cpp
intern/cycles/util/util_types.h
intern/dualcon/intern/Projections.h
intern/dualcon/intern/dualcon_c_api.cpp
intern/elbeem/intern/mvmcoords.h
intern/ghost/intern/GHOST_Context.cpp
intern/ghost/intern/GHOST_DisplayManagerWin32.cpp
intern/ghost/intern/GHOST_SystemPathsWin32.cpp
intern/ghost/intern/GHOST_SystemWin32.cpp
intern/ghost/intern/GHOST_SystemWin32.h
intern/ghost/intern/GHOST_TaskbarWin32.h
intern/guardedalloc/intern/mallocn_intern.h
intern/libmv/ChangeLog
intern/libmv/bundle.sh
intern/libmv/intern/frame_accessor.cc
intern/libmv/intern/frame_accessor.h
intern/libmv/intern/stub.cc
intern/libmv/intern/track_region.cc
intern/libmv/libmv/autotrack/autotrack.cc
intern/libmv/libmv/autotrack/frame_accessor.h
intern/opencolorio/CMakeLists.txt
intern/opensubdiv/opensubdiv_capi.cc
intern/opensubdiv/opensubdiv_capi.h
intern/string/STR_HashedString.h
release/datafiles/blender_icons.svg
release/datafiles/blender_icons16/icon16_outliner_ob_force_field.dat [new file with mode: 0644]
release/datafiles/blender_icons16/icon16_outliner_ob_group_instance.dat [new file with mode: 0644]
release/datafiles/blender_icons32/icon32_outliner_ob_force_field.dat [new file with mode: 0644]
release/datafiles/blender_icons32/icon32_outliner_ob_group_instance.dat [new file with mode: 0644]
release/scripts/modules/bl_i18n_utils/settings.py
release/scripts/modules/bl_i18n_utils/utils_spell_check.py
release/scripts/modules/sys_info.py
release/scripts/startup/bl_operators/object.py
release/scripts/startup/bl_ui/properties_animviz.py
release/scripts/startup/bl_ui/properties_data_armature.py
release/scripts/startup/bl_ui/properties_data_camera.py
release/scripts/startup/bl_ui/properties_data_curve.py
release/scripts/startup/bl_ui/properties_data_lamp.py
release/scripts/startup/bl_ui/properties_data_metaball.py
release/scripts/startup/bl_ui/properties_game.py
release/scripts/startup/bl_ui/properties_grease_pencil_common.py
release/scripts/startup/bl_ui/properties_material.py
release/scripts/startup/bl_ui/properties_object.py
release/scripts/startup/bl_ui/properties_particle.py
release/scripts/startup/bl_ui/properties_physics_dynamicpaint.py
release/scripts/startup/bl_ui/properties_physics_field.py
release/scripts/startup/bl_ui/properties_physics_smoke.py
release/scripts/startup/bl_ui/properties_physics_softbody.py
release/scripts/startup/bl_ui/properties_texture.py
release/scripts/startup/bl_ui/properties_world.py
release/scripts/startup/bl_ui/space_sequencer.py
release/scripts/startup/bl_ui/space_time.py
release/scripts/startup/bl_ui/space_userpref.py
release/scripts/startup/bl_ui/space_view3d.py
release/scripts/startup/bl_ui/space_view3d_toolbar.py
source/blender/alembic/ABC_alembic.h
source/blender/alembic/intern/abc_archive.cc
source/blender/alembic/intern/abc_curves.cc
source/blender/alembic/intern/abc_customdata.cc
source/blender/alembic/intern/abc_exporter.cc
source/blender/alembic/intern/abc_exporter.h
source/blender/alembic/intern/abc_hair.cc
source/blender/alembic/intern/abc_mesh.cc
source/blender/alembic/intern/abc_points.cc
source/blender/alembic/intern/abc_transform.cc
source/blender/alembic/intern/abc_util.cc
source/blender/alembic/intern/alembic_capi.cc
source/blender/blenkernel/BKE_animsys.h
source/blender/blenkernel/BKE_global.h
source/blender/blenkernel/BKE_object.h
source/blender/blenkernel/BKE_texture.h
source/blender/blenkernel/intern/DerivedMesh.c
source/blender/blenkernel/intern/anim_sys.c
source/blender/blenkernel/intern/armature_update.c
source/blender/blenkernel/intern/cdderivedmesh.c
source/blender/blenkernel/intern/constraint.c
source/blender/blenkernel/intern/customdata.c
source/blender/blenkernel/intern/displist.c
source/blender/blenkernel/intern/dynamicpaint.c
source/blender/blenkernel/intern/editderivedmesh.c
source/blender/blenkernel/intern/effect.c
source/blender/blenkernel/intern/group.c
source/blender/blenkernel/intern/image.c
source/blender/blenkernel/intern/mball_tessellate.c
source/blender/blenkernel/intern/mesh_mapping.c
source/blender/blenkernel/intern/mesh_validate.c
source/blender/blenkernel/intern/object_update.c
source/blender/blenkernel/intern/ocean.c
source/blender/blenkernel/intern/packedFile.c
source/blender/blenkernel/intern/particle.c
source/blender/blenkernel/intern/particle_distribute.c
source/blender/blenkernel/intern/particle_system.c
source/blender/blenkernel/intern/pbvh.c
source/blender/blenkernel/intern/rigidbody.c
source/blender/blenkernel/intern/scene.c
source/blender/blenkernel/intern/seqeffects.c
source/blender/blenkernel/intern/seqmodifier.c
source/blender/blenkernel/intern/sequencer.c
source/blender/blenkernel/intern/softbody.c
source/blender/blenkernel/intern/texture.c
source/blender/blenkernel/intern/tracking.c
source/blender/blenkernel/intern/tracking_auto.c
source/blender/blenkernel/intern/tracking_stabilize.c
source/blender/blenkernel/intern/tracking_util.c
source/blender/blenkernel/intern/writeframeserver.c
source/blender/blenkernel/tracking_private.h
source/blender/blenlib/BLI_alloca.h
source/blender/blenlib/BLI_array.h
source/blender/blenlib/BLI_compiler_attrs.h
source/blender/blenlib/BLI_compiler_compat.h
source/blender/blenlib/BLI_fileops.h
source/blender/blenlib/BLI_fileops_types.h
source/blender/blenlib/BLI_ghash.h
source/blender/blenlib/BLI_kdopbvh.h
source/blender/blenlib/BLI_listbase.h
source/blender/blenlib/BLI_math_rotation.h
source/blender/blenlib/BLI_stack.h
source/blender/blenlib/BLI_strict_flags.h
source/blender/blenlib/BLI_sys_types.h
source/blender/blenlib/BLI_task.h
source/blender/blenlib/BLI_winstuff.h
source/blender/blenlib/intern/BLI_ghash.c
source/blender/blenlib/intern/BLI_kdopbvh.c
source/blender/blenlib/intern/array_store.c
source/blender/blenlib/intern/fileops.c
source/blender/blenlib/intern/hash_mm2a.c
source/blender/blenlib/intern/listbase.c
source/blender/blenlib/intern/math_color_blend_inline.c
source/blender/blenlib/intern/math_rotation.c
source/blender/blenlib/intern/path_util.c
source/blender/blenlib/intern/storage.c
source/blender/blenlib/intern/string.c
source/blender/blenlib/intern/string_utf8.c
source/blender/blenlib/intern/task.c
source/blender/blenlib/intern/winstuff.c
source/blender/blenloader/intern/readfile.c
source/blender/blenloader/intern/versioning_260.c
source/blender/blenloader/intern/versioning_270.c
source/blender/blenloader/intern/versioning_legacy.c
source/blender/bmesh/intern/bmesh_opdefines.c
source/blender/bmesh/intern/bmesh_polygon_edgenet.c
source/blender/bmesh/intern/bmesh_queries.c
source/blender/bmesh/operators/bmo_bevel.c
source/blender/bmesh/operators/bmo_removedoubles.c
source/blender/bmesh/tools/bmesh_bevel.c
source/blender/collada/MeshImporter.cpp
source/blender/collada/SkinInfo.cpp
source/blender/compositor/nodes/COM_BrightnessNode.cpp
source/blender/compositor/nodes/COM_ImageNode.cpp
source/blender/compositor/operations/COM_BrightnessOperation.cpp
source/blender/compositor/operations/COM_BrightnessOperation.h
source/blender/compositor/operations/COM_TextureOperation.cpp
source/blender/compositor/operations/COM_TextureOperation.h
source/blender/depsgraph/CMakeLists.txt
source/blender/depsgraph/intern/builder/deg_builder.cc
source/blender/depsgraph/intern/builder/deg_builder_cycle.cc
source/blender/depsgraph/intern/builder/deg_builder_nodes.cc
source/blender/depsgraph/intern/builder/deg_builder_nodes.h
source/blender/depsgraph/intern/builder/deg_builder_nodes_rig.cc
source/blender/depsgraph/intern/builder/deg_builder_nodes_scene.cc
source/blender/depsgraph/intern/builder/deg_builder_pchanmap.cc
source/blender/depsgraph/intern/builder/deg_builder_relations.cc
source/blender/depsgraph/intern/builder/deg_builder_relations.h
source/blender/depsgraph/intern/builder/deg_builder_relations_keys.cc
source/blender/depsgraph/intern/builder/deg_builder_relations_rig.cc
source/blender/depsgraph/intern/builder/deg_builder_relations_scene.cc
source/blender/depsgraph/intern/builder/deg_builder_transitive.cc
source/blender/depsgraph/intern/debug/deg_debug_graphviz.cc
source/blender/depsgraph/intern/depsgraph.cc
source/blender/depsgraph/intern/depsgraph.h
source/blender/depsgraph/intern/depsgraph_build.cc
source/blender/depsgraph/intern/depsgraph_debug.cc
source/blender/depsgraph/intern/depsgraph_eval.cc
source/blender/depsgraph/intern/depsgraph_query.cc
source/blender/depsgraph/intern/depsgraph_tag.cc
source/blender/depsgraph/intern/depsgraph_type_defines.cc
source/blender/depsgraph/intern/depsgraph_types.h
source/blender/depsgraph/intern/eval/deg_eval.cc
source/blender/depsgraph/intern/eval/deg_eval_debug.cc
source/blender/depsgraph/intern/eval/deg_eval_flush.cc
source/blender/depsgraph/intern/nodes/deg_node.cc
source/blender/depsgraph/intern/nodes/deg_node.h
source/blender/depsgraph/intern/nodes/deg_node_component.cc
source/blender/depsgraph/intern/nodes/deg_node_component.h
source/blender/depsgraph/intern/nodes/deg_node_operation.cc
source/blender/depsgraph/intern/nodes/deg_node_operation.h
source/blender/depsgraph/util/deg_util_foreach.h
source/blender/editors/animation/anim_draw.c
source/blender/editors/animation/anim_markers.c
source/blender/editors/animation/keyframes_edit.c
source/blender/editors/armature/pose_slide.c
source/blender/editors/armature/reeb.c
source/blender/editors/curve/editcurve.c
source/blender/editors/gpencil/gpencil_brush.c
source/blender/editors/gpencil/gpencil_edit.c
source/blender/editors/gpencil/gpencil_intern.h
source/blender/editors/gpencil/gpencil_interpolate.c
source/blender/editors/include/UI_icons.h
source/blender/editors/interface/interface_handlers.c
source/blender/editors/io/io_alembic.c
source/blender/editors/mask/mask_ops.c
source/blender/editors/mesh/editmesh_tools.c
source/blender/editors/mesh/meshtools.c
source/blender/editors/object/object_bake_api.c
source/blender/editors/object/object_relations.c
source/blender/editors/physics/dynamicpaint_ops.c
source/blender/editors/physics/particle_edit.c
source/blender/editors/physics/physics_pointcache.c
source/blender/editors/render/render_opengl.c
source/blender/editors/screen/screen_edit.c
source/blender/editors/space_clip/space_clip.c
source/blender/editors/space_clip/tracking_ops.c
source/blender/editors/space_clip/tracking_ops_plane.c
source/blender/editors/space_file/filelist.c
source/blender/editors/space_file/fsmenu.c
source/blender/editors/space_image/image_buttons.c
source/blender/editors/space_nla/nla_buttons.c
source/blender/editors/space_nla/nla_draw.c
source/blender/editors/space_node/drawnode.c
source/blender/editors/space_node/node_add.c
source/blender/editors/space_node/node_view.c
source/blender/editors/space_outliner/outliner_select.c
source/blender/editors/space_outliner/outliner_tools.c
source/blender/editors/space_sequencer/sequencer_edit.c
source/blender/editors/space_sequencer/sequencer_select.c
source/blender/editors/space_text/CMakeLists.txt
source/blender/editors/space_text/space_text.c
source/blender/editors/space_text/text_autocomplete.c
source/blender/editors/space_text/text_format.h
source/blender/editors/space_text/text_format_pov.c [new file with mode: 0644]
source/blender/editors/space_text/text_format_pov_ini.c [new file with mode: 0644]
source/blender/editors/space_view3d/drawobject.c
source/blender/editors/space_view3d/view3d_draw.c
source/blender/editors/space_view3d/view3d_iterators.c
source/blender/editors/space_view3d/view3d_view.c
source/blender/editors/transform/transform.c
source/blender/editors/transform/transform_conversions.c
source/blender/editors/transform/transform_generics.c
source/blender/editors/transform/transform_input.c
source/blender/editors/transform/transform_manipulator.c
source/blender/editors/transform/transform_orientations.c
source/blender/editors/util/numinput.c
source/blender/editors/uvedit/uvedit_parametrizer.c
source/blender/editors/uvedit/uvedit_smart_stitch.c
source/blender/freestyle/intern/blender_interface/BlenderStrokeRenderer.cpp
source/blender/gpu/intern/gpu_debug.c
source/blender/gpu/shaders/gpu_shader_material.glsl
source/blender/imbuf/intern/IMB_anim.h
source/blender/imbuf/intern/anim_movie.c
source/blender/imbuf/intern/cineon/dpxlib.c
source/blender/imbuf/intern/cineon/logImageCore.c
source/blender/imbuf/intern/colormanagement.c
source/blender/imbuf/intern/dds/dds_api.cpp
source/blender/imbuf/intern/filter.c
source/blender/imbuf/intern/oiio/openimageio_api.cpp
source/blender/imbuf/intern/openexr/openexr_api.cpp
source/blender/imbuf/intern/png.c
source/blender/imbuf/intern/radiance_hdr.c
source/blender/makesrna/intern/makesrna.c
source/blender/makesrna/intern/rna_access.c
source/blender/makesrna/intern/rna_define.c
source/blender/makesrna/intern/rna_fluidsim.c
source/blender/makesrna/intern/rna_image.c
source/blender/makesrna/intern/rna_nodetree.c
source/blender/makesrna/intern/rna_object.c
source/blender/makesrna/intern/rna_particle.c
source/blender/makesrna/intern/rna_scene.c
source/blender/makesrna/intern/rna_scene_api.c
source/blender/makesrna/intern/rna_sequencer.c
source/blender/makesrna/intern/rna_space.c
source/blender/makesrna/intern/rna_userdef.c
source/blender/makesrna/intern/rna_wm.c
source/blender/modifiers/intern/MOD_array.c
source/blender/modifiers/intern/MOD_displace.c
source/blender/modifiers/intern/MOD_meshcache_mdd.c
source/blender/modifiers/intern/MOD_meshcache_pc2.c
source/blender/modifiers/intern/MOD_weightvgedit.c
source/blender/modifiers/intern/MOD_weightvgmix.c
source/blender/nodes/NOD_static_types.h
source/blender/nodes/composite/nodes/node_composite_brightness.c
source/blender/nodes/composite/nodes/node_composite_image.c
source/blender/nodes/intern/node_exec.c
source/blender/nodes/shader/node_shader_util.c
source/blender/nodes/shader/nodes/node_shader_bsdf_principled.c
source/blender/nodes/shader/nodes/node_shader_normal_map.c
source/blender/python/intern/CMakeLists.txt
source/blender/python/intern/bpy_app.c
source/blender/python/intern/bpy_app_handlers.c
source/blender/python/intern/bpy_app_opensubdiv.c [new file with mode: 0644]
source/blender/python/intern/bpy_app_opensubdiv.h [moved from source/blender/depsgraph/util/deg_util_hash.h with 61% similarity]
source/blender/python/intern/bpy_app_sdl.c
source/blender/python/intern/bpy_interface.c
source/blender/python/intern/bpy_operator_wrap.c
source/blender/python/intern/bpy_operator_wrap.h
source/blender/python/intern/bpy_rna.c
source/blender/python/mathutils/mathutils_Color.c
source/blender/python/mathutils/mathutils_Euler.c
source/blender/python/mathutils/mathutils_Matrix.c
source/blender/python/mathutils/mathutils_Quaternion.c
source/blender/render/intern/source/bake_api.c
source/blender/render/intern/source/convertblender.c
source/blender/render/intern/source/render_texture.c
source/blender/render/intern/source/rendercore.c
source/blender/render/intern/source/sunsky.c
source/blender/render/intern/source/volumetric.c
source/blender/render/intern/source/zbuf.c
source/blender/windowmanager/intern/wm_files.c
source/blenderplayer/CMakeLists.txt
source/blenderplayer/bad_level_call_stubs/stubs.c
source/creator/CMakeLists.txt
source/gameengine/Expressions/intern/HashedPtr.cpp
source/gameengine/Expressions/intern/InputParser.cpp
source/gameengine/Expressions/intern/PyObjectPlus.cpp
source/gameengine/GameLogic/SCA_PropertySensor.cpp
source/gameengine/Ketsji/KX_IPO_SGController.cpp
source/gameengine/Ketsji/KX_KetsjiEngine.cpp
source/gameengine/Ketsji/KX_ObstacleSimulation.cpp
source/gameengine/Ketsji/KX_SoundActuator.cpp
source/gameengine/VideoTexture/ImageBase.cpp
tests/gtests/alembic/CMakeLists.txt
tests/gtests/alembic/abc_export_test.cc [new file with mode: 0644]
tests/gtests/blenlib/BLI_kdopbvh_test.cc [new file with mode: 0644]
tests/gtests/blenlib/CMakeLists.txt
tests/python/CMakeLists.txt
tests/python/bl_alembic_import_test.py
tests/python/bl_keymap_completeness.py
tests/python/pep8.py
tests/python/rna_info_dump.py

index 1dac082459f17bd2ec3e5e44d0dd39e6a3c04d6d..66ee215864dab3d15b2a277ac7be85f791dd85ce 100644 (file)
@@ -229,7 +229,7 @@ mark_as_advanced(BUILDINFO_OVERRIDE_TIME)
 
 option(WITH_IK_ITASC      "Enable ITASC IK solver (only disable for development & for incompatible C++ compilers)" ON)
 option(WITH_IK_SOLVER     "Enable Legacy IK solver (only disable for development)" ON)
-option(WITH_FFTW3         "Enable FFTW3 support (Used for smoke and audio effects)" ${_init_FFTW3})
+option(WITH_FFTW3         "Enable FFTW3 support (Used for smoke, ocean sim, and audio effects)" ${_init_FFTW3})
 option(WITH_BULLET        "Enable Bullet (Physics Engine)" ON)
 option(WITH_SYSTEM_BULLET "Use the systems bullet library (currently unsupported due to missing features in upstream!)" )
 mark_as_advanced(WITH_SYSTEM_BULLET)
@@ -1380,10 +1380,20 @@ if(CMAKE_COMPILER_IS_GNUCC)
                ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_MISSING_DECLARATIONS -Wmissing-declarations)
        endif()
 
+       # Use 'ATTR_FALLTHROUGH' macro to suppress.
+       if(CMAKE_COMPILER_IS_GNUCC AND (NOT "${CMAKE_C_COMPILER_VERSION}" VERSION_LESS "7.0"))
+               ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_IMPLICIT_FALLTHROUGH -Wimplicit-fallthrough=5)
+               ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_IMPLICIT_FALLTHROUGH -Wimplicit-fallthrough=5)
+       endif()
+
        # flags to undo strict flags
        ADD_CHECK_C_COMPILER_FLAG(CC_REMOVE_STRICT_FLAGS C_WARN_NO_DEPRECATED_DECLARATIONS -Wno-deprecated-declarations)
        ADD_CHECK_C_COMPILER_FLAG(CC_REMOVE_STRICT_FLAGS C_WARN_NO_UNUSED_PARAMETER        -Wno-unused-parameter)
 
+       if(CMAKE_COMPILER_IS_GNUCC AND (NOT "${CMAKE_C_COMPILER_VERSION}" VERSION_LESS "7.0"))
+               ADD_CHECK_C_COMPILER_FLAG(CC_REMOVE_STRICT_FLAGS C_WARN_NO_IMPLICIT_FALLTHROUGH    -Wno-implicit-fallthrough)
+       endif()
+
        if(NOT APPLE)
                ADD_CHECK_C_COMPILER_FLAG(CC_REMOVE_STRICT_FLAGS C_WARN_NO_ERROR_UNUSED_BUT_SET_VARIABLE -Wno-error=unused-but-set-variable)
        endif()
index 6b85ad3dde1896aebb8934812b25db4e583f8a0c..9e256a78a598513a505f548f4d8c1bad57fc70bf 100755 (executable)
@@ -289,7 +289,7 @@ NO_BUILD=false
 NO_CONFIRM=false
 USE_CXX11=false
 
-PYTHON_VERSION="3.5.2"
+PYTHON_VERSION="3.5.3"
 PYTHON_VERSION_MIN="3.5"
 PYTHON_FORCE_BUILD=false
 PYTHON_FORCE_REBUILD=false
@@ -322,8 +322,8 @@ OPENEXR_FORCE_REBUILD=false
 OPENEXR_SKIP=false
 _with_built_openexr=false
 
-OIIO_VERSION="1.7.13"
-OIIO_VERSION_MIN="1.7.13"
+OIIO_VERSION="1.7.15"
+OIIO_VERSION_MIN="1.7.15"
 OIIO_VERSION_MAX="1.9.0"  # UNKNOWN currently # Not supported by current OSL...
 OIIO_FORCE_BUILD=false
 OIIO_FORCE_REBUILD=false
@@ -366,8 +366,7 @@ ALEMBIC_FORCE_BUILD=false
 ALEMBIC_FORCE_REBUILD=false
 ALEMBIC_SKIP=false
 
-# Version??
-OPENCOLLADA_VERSION="1.3"
+OPENCOLLADA_VERSION="1.6.51"
 OPENCOLLADA_FORCE_BUILD=false
 OPENCOLLADA_FORCE_REBUILD=false
 OPENCOLLADA_SKIP=false
@@ -737,7 +736,10 @@ _boost_version_nodots=`echo "$BOOST_VERSION" | sed -r 's/\./_/g'`
 BOOST_SOURCE=( "http://sourceforge.net/projects/boost/files/boost/$BOOST_VERSION/boost_$_boost_version_nodots.tar.bz2/download" )
 BOOST_BUILD_MODULES="--with-system --with-filesystem --with-thread --with-regex --with-locale --with-date_time --with-wave --with-iostreams --with-python --with-program_options"
 
+OCIO_USE_REPO=true
 OCIO_SOURCE=( "https://github.com/imageworks/OpenColorIO/tarball/v$OCIO_VERSION" )
+OCIO_SOURCE_REPO=( "https://github.com/imageworks/OpenColorIO.git" )
+OCIO_SOURCE_REPO_UID="6de971097c7f552300f669ed69ca0b6cf5a70843"
 
 OPENEXR_USE_REPO=false
 OPENEXR_SOURCE=( "http://download.savannah.nongnu.org/releases/openexr/openexr-$OPENEXR_VERSION.tar.gz" )
@@ -786,7 +788,7 @@ ALEMBIC_SOURCE=( "https://github.com/alembic/alembic/archive/${ALEMBIC_VERSION}.
 # ALEMBIC_SOURCE_REPO_BRANCH="master"
 
 OPENCOLLADA_SOURCE=( "https://github.com/KhronosGroup/OpenCOLLADA.git" )
-OPENCOLLADA_REPO_UID="3335ac164e68b2512a40914b14c74db260e6ff7d"
+OPENCOLLADA_REPO_UID="0c2cdc17c22cf42050e4d42154bed2176363549c"
 OPENCOLLADA_REPO_BRANCH="master"
 
 FFMPEG_SOURCE=( "http://ffmpeg.org/releases/ffmpeg-$FFMPEG_VERSION.tar.bz2" )
@@ -1268,7 +1270,7 @@ compile_OCIO() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  ocio_magic=1
+  ocio_magic=2
   _init_ocio
 
   # Clean install if needed!
@@ -1285,14 +1287,27 @@ compile_OCIO() {
     if [ ! -d $_src ]; then
       INFO "Downloading OpenColorIO-$OCIO_VERSION"
       mkdir -p $SRC
-      download OCIO_SOURCE[@] $_src.tar.gz
 
-      INFO "Unpacking OpenColorIO-$OCIO_VERSION"
-      tar -C $SRC --transform "s,(.*/?)imageworks-OpenColorIO[^/]*(.*),\1OpenColorIO-$OCIO_VERSION\2,x" \
-          -xf $_src.tar.gz
+      if [ "$OCIO_USE_REPO" = true ]; then
+        git clone ${OCIO_SOURCE_REPO[0]} $_src
+      else
+        download OCIO_SOURCE[@] $_src.tar.gz
+        INFO "Unpacking OpenColorIO-$OCIO_VERSION"
+        tar -C $SRC --transform "s,(.*/?)imageworks-OpenColorIO[^/]*(.*),\1OpenColorIO-$OCIO_VERSION\2,x" \
+            -xf $_src.tar.gz
+      fi
+
     fi
 
     cd $_src
+
+    if [ "$OCIO_USE_REPO" = true ]; then
+      # XXX For now, always update from latest repo...
+      git pull origin master
+      git checkout $OCIO_SOURCE_REPO_UID
+      git reset --hard
+    fi
+
     # Always refresh the whole build!
     if [ -d build ]; then
       rm -rf build
@@ -1498,7 +1513,6 @@ compile_OPENEXR() {
     if [ "$OPENEXR_USE_REPO" = true ]; then
       # XXX For now, always update from latest repo...
       git pull origin master
-      # Stick to same rev as windows' libs...
       git checkout $OPENEXR_SOURCE_REPO_UID
       git reset --hard
       oiio_src_path="../OpenEXR"
index 67b628f668a24184bab73fab95eb55e03eec9b91..157e244e210135a7d69e8c9dc0b25065382f46b6 100644 (file)
@@ -49,7 +49,7 @@ def get_platform(filename):
     tokens = filename.split("-")
     platforms = ('osx', 'mac', 'bsd',
                  'win', 'linux', 'source',
-                 'irix', 'solaris', 'mingw')
+                 'irix', 'solaris')
     platform_tokens = []
     found = False
 
index 1c10d5a71debf5d92e40d3ecdf2cfedcb3d2f4fd..63bc520ea15b5d03d023bacf3df3285d5da17327 100644 (file)
@@ -83,6 +83,7 @@ FOREACH(COMPONENT ${_opencollada_FIND_INCLUDES})
       # but this is less trouble, just looks strange.
       include/opencollada/${COMPONENT}
       include/${COMPONENT}/include
+      include/${COMPONENT}
     HINTS
       ${_opencollada_SEARCH_DIRS}
     )
index ecde50ff156301cf72a63ee5d25334f0b4b368d1..62e2ce1636d466e45727d3c07138ca31ab043229 100644 (file)
@@ -63,17 +63,10 @@ if(UNIX AND NOT APPLE)
        set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
 elseif(WIN32)
        set(WITH_JACK                OFF CACHE BOOL "" FORCE)
-       if(NOT CMAKE_COMPILER_IS_GNUCC)
-               set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
-       else()
-               # MinGW exceptions
-               set(WITH_OPENSUBDIV          OFF CACHE BOOL "" FORCE)
-               set(WITH_CODEC_SNDFILE       OFF CACHE BOOL "" FORCE)
-               set(WITH_CYCLES_OSL          OFF CACHE BOOL "" FORCE)
-       endif()
+       set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
 elseif(APPLE)
        set(WITH_JACK                ON  CACHE BOOL "" FORCE)
-       set(WITH_CODEC_QUICKTIME     OFF  CACHE BOOL "" FORCE)
+       set(WITH_CODEC_QUICKTIME     OFF CACHE BOOL "" FORCE)
        set(WITH_OPENSUBDIV          OFF CACHE BOOL "" FORCE)
 
 #      include("${CMAKE_CURRENT_SOURCE_DIR}/../platform/platform_apple_xcode.cmake")
index f7239559fb8c50a02869808997d5a925bae42d76..3e2e26e6a4433c992e1c15061703e1cef6f38abe 100644 (file)
@@ -64,17 +64,10 @@ if(UNIX AND NOT APPLE)
        set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
 elseif(WIN32)
        set(WITH_JACK                OFF CACHE BOOL "" FORCE)
-       if(NOT CMAKE_COMPILER_IS_GNUCC)
-               set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
-       else()
-               # MinGW exceptions
-               set(WITH_OPENSUBDIV          OFF CACHE BOOL "" FORCE)
-               set(WITH_CODEC_SNDFILE       OFF CACHE BOOL "" FORCE)
-               set(WITH_CYCLES_OSL          OFF CACHE BOOL "" FORCE)
-       endif()
+       set(WITH_OPENSUBDIV          ON  CACHE BOOL "" FORCE)
 elseif(APPLE)
        set(WITH_JACK                ON  CACHE BOOL "" FORCE)
-       set(WITH_CODEC_QUICKTIME     OFF  CACHE BOOL "" FORCE)
+       set(WITH_CODEC_QUICKTIME     OFF CACHE BOOL "" FORCE)
        set(WITH_OPENSUBDIV          OFF CACHE BOOL "" FORCE)
 
 #      include("${CMAKE_CURRENT_SOURCE_DIR}/../platform/platform_apple_xcode.cmake")
index f62e55941a7a95c38b5a47af16804951af274112..6303119773d1a74bc6b52c5bad214005a17085fb 100644 (file)
@@ -1247,17 +1247,6 @@ endfunction()
 # hacks to override initial project settings
 # these macros must be called directly before/after project(Blender)
 macro(blender_project_hack_pre)
-       # ----------------
-       # MINGW HACK START
-       # ignore system set flag, use our own
-       # must be before project(...)
-       # if the user wants to add their own its ok after first run.
-       if(DEFINED CMAKE_C_STANDARD_LIBRARIES)
-               set(_reset_standard_libraries OFF)
-       else()
-               set(_reset_standard_libraries ON)
-       endif()
-
        # ------------------
        # GCC -O3 HACK START
        # needed because O3 can cause problems but
@@ -1276,25 +1265,6 @@ endmacro()
 
 
 macro(blender_project_hack_post)
-       # --------------
-       # MINGW HACK END
-       if(_reset_standard_libraries)
-               # Must come after projecINCt(...)
-               #
-               # MINGW workaround for -ladvapi32 being included which surprisingly causes
-               # string formatting of floats, eg: printf("%.*f", 3, value). to crash blender
-               # with a meaningless stack trace. by overriding this flag we ensure we only
-               # have libs we define.
-               set(CMAKE_C_STANDARD_LIBRARIES "" CACHE STRING "" FORCE)
-               set(CMAKE_CXX_STANDARD_LIBRARIES "" CACHE STRING "" FORCE)
-               mark_as_advanced(
-                       CMAKE_C_STANDARD_LIBRARIES
-                       CMAKE_CXX_STANDARD_LIBRARIES
-               )
-       endif()
-       unset(_reset_standard_libraries)
-
-
        # ----------------
        # GCC -O3 HACK END
        if(_reset_standard_cflags_rel)
index 631973b758b01332aeeb99235c3427d446b5a911..2f5d41dac32bcc529ae7fa7d0207c532ce0cdb56 100644 (file)
@@ -27,12 +27,10 @@ add_definitions(-DWIN32)
 
 if(MSVC)
        include(platform_win32_msvc)
-elseif(CMAKE_COMPILER_IS_GNUCC)
-       include(platform_win32_mingw)
+else()
+       message(FATAL_ERROR "Compiler is unsupported")
 endif()
 
-# Things common to both mingw and MSVC  should go here
-
 set(WINTAB_INC ${LIBDIR}/wintab/include)
 
 if(WITH_OPENAL)
diff --git a/build_files/cmake/platform/platform_win32_mingw.cmake b/build_files/cmake/platform/platform_win32_mingw.cmake
deleted file mode 100644 (file)
index 216568b..0000000
+++ /dev/null
@@ -1,302 +0,0 @@
-# ***** BEGIN GPL LICENSE BLOCK *****
-#
-# This program is free software; you can redistribute it and/or
-# modify it under the terms of the GNU General Public License
-# as published by the Free Software Foundation; either version 2
-# of the License, or (at your option) any later version.
-#
-# This program is distributed in the hope that it will be useful,
-# but WITHOUT ANY WARRANTY; without even the implied warranty of
-# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-# GNU General Public License for more details.
-#
-# You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software Foundation,
-# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-#
-# The Original Code is Copyright (C) 2016, Blender Foundation
-# All rights reserved.
-#
-# Contributor(s): Sergey Sharybin.
-#
-# ***** END GPL LICENSE BLOCK *****
-
-# Libraries configuration for Windows when compiling with MinGW.
-
-# keep GCC specific stuff here
-include(CheckCSourceCompiles)
-# Setup 64bit and 64bit windows systems
-CHECK_C_SOURCE_COMPILES("
-       #ifndef __MINGW64__
-       #error
-       #endif
-       int main(void) { return 0; }
-       "
-       WITH_MINGW64
-)
-
-if(NOT DEFINED LIBDIR)
-       if(WITH_MINGW64)
-               message(STATUS "Compiling for 64 bit with MinGW-w64.")
-               set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/mingw64)
-       else()
-               message(STATUS "Compiling for 32 bit with MinGW-w32.")
-               set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/mingw32)
-
-               if(WITH_RAYOPTIMIZATION)
-                       message(WARNING "MinGW-w32 is known to be unstable with 'WITH_RAYOPTIMIZATION' option enabled.")
-               endif()
-       endif()
-else()
-       message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
-endif()
-if(NOT EXISTS "${LIBDIR}/")
-       message(FATAL_ERROR "Windows requires pre-compiled libs at: '${LIBDIR}'")
-endif()
-
-list(APPEND PLATFORM_LINKLIBS
-       -lshell32 -lshfolder -lgdi32 -lmsvcrt -lwinmm -lmingw32 -lm -lws2_32
-       -lz -lstdc++ -lole32 -luuid -lwsock32 -lpsapi -ldbghelp
-)
-
-if(WITH_INPUT_IME)
-       list(APPEND PLATFORM_LINKLIBS -limm32)
-endif()
-
-set(PLATFORM_CFLAGS "-pipe -funsigned-char -fno-strict-aliasing")
-
-if(WITH_MINGW64)
-       set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fpermissive")
-       list(APPEND PLATFORM_LINKLIBS -lpthread)
-
-       add_definitions(-DFREE_WINDOWS64 -DMS_WIN64)
-endif()
-
-add_definitions(-D_LARGEFILE_SOURCE -D_FILE_OFFSET_BITS=64 -D_LARGEFILE64_SOURCE)
-
-add_definitions(-DFREE_WINDOWS)
-
-set(PNG "${LIBDIR}/png")
-set(PNG_INCLUDE_DIRS "${PNG}/include")
-set(PNG_LIBPATH ${PNG}/lib) # not cmake defined
-
-if(WITH_MINGW64)
-       set(JPEG_LIBRARIES jpeg)
-else()
-       set(JPEG_LIBRARIES libjpeg)
-endif()
-set(PNG_LIBRARIES png)
-
-set(ZLIB ${LIBDIR}/zlib)
-set(ZLIB_INCLUDE_DIRS ${ZLIB}/include)
-set(ZLIB_LIBPATH ${ZLIB}/lib)
-set(ZLIB_LIBRARIES z)
-
-set(JPEG "${LIBDIR}/jpeg")
-set(JPEG_INCLUDE_DIR "${JPEG}/include")
-set(JPEG_LIBPATH ${JPEG}/lib) # not cmake defined
-
-# comes with own pthread library
-if(NOT WITH_MINGW64)
-       set(PTHREADS ${LIBDIR}/pthreads)
-       #set(PTHREADS_INCLUDE_DIRS ${PTHREADS}/include)
-       set(PTHREADS_LIBPATH ${PTHREADS}/lib)
-       set(PTHREADS_LIBRARIES pthreadGC2)
-endif()
-
-set(FREETYPE ${LIBDIR}/freetype)
-set(FREETYPE_INCLUDE_DIRS ${FREETYPE}/include ${FREETYPE}/include/freetype2)
-set(FREETYPE_LIBPATH ${FREETYPE}/lib)
-set(FREETYPE_LIBRARY freetype)
-
-if(WITH_FFTW3)
-       set(FFTW3 ${LIBDIR}/fftw3)
-       set(FFTW3_LIBRARIES fftw3)
-       set(FFTW3_INCLUDE_DIRS ${FFTW3}/include)
-       set(FFTW3_LIBPATH ${FFTW3}/lib)
-endif()
-
-if(WITH_OPENCOLLADA)
-       set(OPENCOLLADA ${LIBDIR}/opencollada)
-       set(OPENCOLLADA_INCLUDE_DIRS
-               ${OPENCOLLADA}/include/opencollada/COLLADAStreamWriter
-               ${OPENCOLLADA}/include/opencollada/COLLADABaseUtils
-               ${OPENCOLLADA}/include/opencollada/COLLADAFramework
-               ${OPENCOLLADA}/include/opencollada/COLLADASaxFrameworkLoader
-               ${OPENCOLLADA}/include/opencollada/GeneratedSaxParser
-       )
-       set(OPENCOLLADA_LIBPATH ${OPENCOLLADA}/lib/opencollada)
-       set(OPENCOLLADA_LIBRARIES
-               OpenCOLLADAStreamWriter
-               OpenCOLLADASaxFrameworkLoader
-               OpenCOLLADAFramework
-               OpenCOLLADABaseUtils
-               GeneratedSaxParser
-               UTF MathMLSolver buffer ftoa xml
-       )
-       set(PCRE_LIBRARIES pcre)
-endif()
-
-if(WITH_CODEC_FFMPEG)
-       set(FFMPEG ${LIBDIR}/ffmpeg)
-       set(FFMPEG_INCLUDE_DIRS ${FFMPEG}/include)
-       if(WITH_MINGW64)
-               set(FFMPEG_LIBRARIES avcodec.dll avformat.dll avdevice.dll avutil.dll swscale.dll swresample.dll)
-       else()
-               set(FFMPEG_LIBRARIES avcodec-55 avformat-55 avdevice-55 avutil-52 swscale-2)
-       endif()
-       set(FFMPEG_LIBPATH ${FFMPEG}/lib)
-endif()
-
-if(WITH_IMAGE_OPENEXR)
-       set(OPENEXR ${LIBDIR}/openexr)
-       set(OPENEXR_INCLUDE_DIR ${OPENEXR}/include)
-       set(OPENEXR_INCLUDE_DIRS ${OPENEXR}/include/OpenEXR)
-       set(OPENEXR_LIBRARIES Half IlmImf Imath IlmThread Iex)
-       set(OPENEXR_LIBPATH ${OPENEXR}/lib)
-endif()
-
-if(WITH_IMAGE_TIFF)
-       set(TIFF ${LIBDIR}/tiff)
-       set(TIFF_LIBRARY tiff)
-       set(TIFF_INCLUDE_DIR ${TIFF}/include)
-       set(TIFF_LIBPATH ${TIFF}/lib)
-endif()
-
-if(WITH_JACK)
-       set(JACK ${LIBDIR}/jack)
-       set(JACK_INCLUDE_DIRS ${JACK}/include/jack ${JACK}/include)
-       set(JACK_LIBRARIES jack)
-       set(JACK_LIBPATH ${JACK}/lib)
-
-       # TODO, gives linking errors, force off
-       set(WITH_JACK OFF)
-endif()
-
-if(WITH_PYTHON)
-       # normally cached but not since we include them with blender
-       set(PYTHON_VERSION 3.5) #  CACHE STRING)
-       string(REPLACE "." "" _PYTHON_VERSION_NO_DOTS ${PYTHON_VERSION})
-       set(PYTHON_INCLUDE_DIR "${LIBDIR}/python/include/python${PYTHON_VERSION}")  # CACHE PATH)
-       set(PYTHON_LIBRARY "${LIBDIR}/python/lib/python${_PYTHON_VERSION_NO_DOTS}mw.lib")  # CACHE FILEPATH)
-       unset(_PYTHON_VERSION_NO_DOTS)
-
-       # uncached vars
-       set(PYTHON_INCLUDE_DIRS "${PYTHON_INCLUDE_DIR}")
-       set(PYTHON_LIBRARIES  "${PYTHON_LIBRARY}")
-endif()
-
-if(WITH_BOOST)
-       set(BOOST ${LIBDIR}/boost)
-       set(BOOST_INCLUDE_DIR ${BOOST}/include)
-       if(WITH_MINGW64)
-               set(BOOST_POSTFIX "mgw47-mt-s-1_49")
-               set(BOOST_DEBUG_POSTFIX "mgw47-mt-sd-1_49")
-       else()
-               set(BOOST_POSTFIX "mgw46-mt-s-1_49")
-               set(BOOST_DEBUG_POSTFIX "mgw46-mt-sd-1_49")
-       endif()
-       set(BOOST_LIBRARIES
-               optimized boost_date_time-${BOOST_POSTFIX} boost_filesystem-${BOOST_POSTFIX}
-               boost_regex-${BOOST_POSTFIX}
-               boost_system-${BOOST_POSTFIX} boost_thread-${BOOST_POSTFIX}
-               debug boost_date_time-${BOOST_DEBUG_POSTFIX} boost_filesystem-${BOOST_DEBUG_POSTFIX}
-               boost_regex-${BOOST_DEBUG_POSTFIX}
-               boost_system-${BOOST_DEBUG_POSTFIX} boost_thread-${BOOST_DEBUG_POSTFIX})
-       if(WITH_INTERNATIONAL)
-               set(BOOST_LIBRARIES ${BOOST_LIBRARIES}
-                       optimized boost_locale-${BOOST_POSTFIX}
-                       debug boost_locale-${BOOST_DEBUG_POSTFIX}
-               )
-       endif()
-       if(WITH_CYCLES_OSL)
-               set(BOOST_LIBRARIES ${BOOST_LIBRARIES}
-                       optimized boost_wave-${BOOST_POSTFIX}
-                       debug boost_wave-${BOOST_DEBUG_POSTFIX}
-               )
-       endif()
-       set(BOOST_LIBPATH ${BOOST}/lib)
-       set(BOOST_DEFINITIONS "-DBOOST_ALL_NO_LIB -DBOOST_THREAD_USE_LIB ")
-endif()
-
-if(WITH_OPENIMAGEIO)
-       set(OPENIMAGEIO ${LIBDIR}/openimageio)
-       set(OPENIMAGEIO_INCLUDE_DIRS ${OPENIMAGEIO}/include)
-       set(OPENIMAGEIO_LIBRARIES OpenImageIO)
-       set(OPENIMAGEIO_LIBPATH ${OPENIMAGEIO}/lib)
-       set(OPENIMAGEIO_DEFINITIONS "")
-       set(OPENIMAGEIO_IDIFF "${OPENIMAGEIO}/bin/idiff.exe")
-endif()
-
-if(WITH_LLVM)
-       set(LLVM_ROOT_DIR ${LIBDIR}/llvm CACHE PATH     "Path to the LLVM installation")
-       set(LLVM_LIBPATH ${LLVM_ROOT_DIR}/lib)
-       # Explicitly set llvm lib order.
-       #---- WARNING ON GCC ORDER OF LIBS IS IMPORTANT, DO NOT CHANGE! ---------
-       set(LLVM_LIBRARY LLVMSelectionDAG LLVMCodeGen LLVMScalarOpts LLVMAnalysis LLVMArchive
-               LLVMAsmParser LLVMAsmPrinter
-               LLVMBitReader LLVMBitWriter
-               LLVMDebugInfo LLVMExecutionEngine
-               LLVMInstCombine LLVMInstrumentation
-               LLVMInterpreter LLVMJIT
-               LLVMLinker LLVMMC
-               LLVMMCDisassembler LLVMMCJIT
-               LLVMMCParser LLVMObject
-               LLVMRuntimeDyld
-               LLVMSupport
-               LLVMTableGen LLVMTarget
-               LLVMTransformUtils LLVMVectorize
-               LLVMX86AsmParser LLVMX86AsmPrinter
-               LLVMX86CodeGen LLVMX86Desc
-               LLVMX86Disassembler LLVMX86Info
-               LLVMX86Utils LLVMipa
-               LLVMipo LLVMCore)
-       # imagehelp is needed by LLVM 3.1 on MinGW, check lib\Support\Windows\Signals.inc
-       list(APPEND PLATFORM_LINKLIBS -limagehlp)
-endif()
-
-if(WITH_OPENCOLORIO)
-       set(OPENCOLORIO ${LIBDIR}/opencolorio)
-       set(OPENCOLORIO_INCLUDE_DIRS ${OPENCOLORIO}/include)
-       set(OPENCOLORIO_LIBRARIES OpenColorIO)
-       set(OPENCOLORIO_LIBPATH ${OPENCOLORIO}/lib)
-       set(OPENCOLORIO_DEFINITIONS)
-endif()
-
-if(WITH_SDL)
-       set(SDL ${LIBDIR}/sdl)
-       set(SDL_INCLUDE_DIR ${SDL}/include)
-       set(SDL_LIBRARY SDL)
-       set(SDL_LIBPATH ${SDL}/lib)
-endif()
-
-if(WITH_OPENVDB)
-       set(OPENVDB ${LIBDIR}/openvdb)
-       set(OPENVDB_INCLUDE_DIRS ${OPENVDB}/include)
-       set(OPENVDB_LIBRARIES openvdb ${TBB_LIBRARIES})
-       set(OPENVDB_LIBPATH ${LIBDIR}/openvdb/lib)
-       set(OPENVDB_DEFINITIONS)
-endif()
-
-if(WITH_ALEMBIC)
-       # TODO(sergey): For until someone drops by and compiles libraries for
-       # MinGW we allow users to compile their own Alembic library and use
-       # that via find_package(),
-       #
-       # Once precompiled libraries are there we'll use hardcoded locations.
-       find_package_wrapper(Alembic)
-       if(WITH_ALEMBIC_HDF5)
-               set(HDF5_ROOT_DIR ${LIBDIR}/hdf5)
-               find_package_wrapper(HDF5)
-       endif()
-       if(NOT ALEMBIC_FOUND OR (WITH_ALEMBIC_HDF5 AND NOT HDF5_FOUND))
-               set(WITH_ALEMBIC OFF)
-               set(WITH_ALEMBIC_HDF5 OFF)
-       endif()
-endif()
-
-set(PLATFORM_LINKFLAGS "-Xlinker --stack=2097152")
-
-## DISABLE - causes linking errors
-## for re-distribution, so users dont need mingw installed
-# set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} -static-libgcc -static-libstdc++")
index 6447559f62cfd861ac91f838876b5153134f28b7..3b417f79cbe04c933b4098054e608d8727c614aa 100644 (file)
@@ -116,7 +116,6 @@ set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} /NODEFAULTLIB:msvcrt.lib /NODEFAUL
 set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} /ignore:4049 /ignore:4217 /ignore:4221")
 set(CMAKE_STATIC_LINKER_FLAGS "${CMAKE_STATIC_LINKER_FLAGS} /ignore:4221")
 
-# MSVC only, Mingw doesnt need
 if(CMAKE_CL_64)
        set(PLATFORM_LINKFLAGS "/MACHINE:X64 ${PLATFORM_LINKFLAGS}")
 else()
@@ -453,32 +452,28 @@ if(WITH_MOD_CLOTH_ELTOPO)
 endif()
 
 if(WITH_OPENSUBDIV OR WITH_CYCLES_OPENSUBDIV)
-    set(OPENSUBDIV_INCLUDE_DIR ${LIBDIR}/opensubdiv/include)
-    set(OPENSUBDIV_LIBPATH ${LIBDIR}/opensubdiv/lib)
-    set(OPENSUBDIV_LIBRARIES    optimized ${OPENSUBDIV_LIBPATH}/osdCPU.lib 
-                                optimized ${OPENSUBDIV_LIBPATH}/osdGPU.lib
-                                debug ${OPENSUBDIV_LIBPATH}/osdCPU_d.lib 
-                                debug ${OPENSUBDIV_LIBPATH}/osdGPU_d.lib
-                                )
-    set(OPENSUBDIV_HAS_OPENMP TRUE)
+       set(OPENSUBDIV_INCLUDE_DIR ${LIBDIR}/opensubdiv/include)
+       set(OPENSUBDIV_LIBPATH ${LIBDIR}/opensubdiv/lib)
+       set(OPENSUBDIV_LIBRARIES
+               optimized ${OPENSUBDIV_LIBPATH}/osdCPU.lib
+               optimized ${OPENSUBDIV_LIBPATH}/osdGPU.lib
+               debug ${OPENSUBDIV_LIBPATH}/osdCPU_d.lib
+               debug ${OPENSUBDIV_LIBPATH}/osdGPU_d.lib
+       )
+       set(OPENSUBDIV_HAS_OPENMP TRUE)
        set(OPENSUBDIV_HAS_TBB FALSE)
        set(OPENSUBDIV_HAS_OPENCL TRUE)
        set(OPENSUBDIV_HAS_CUDA FALSE)
        set(OPENSUBDIV_HAS_GLSL_TRANSFORM_FEEDBACK TRUE)
        set(OPENSUBDIV_HAS_GLSL_COMPUTE TRUE)
-    windows_find_package(OpenSubdiv)
+       windows_find_package(OpenSubdiv)
 endif()
 
 if(WITH_SDL)
        set(SDL ${LIBDIR}/sdl)
        set(SDL_INCLUDE_DIR ${SDL}/include)
        set(SDL_LIBPATH ${SDL}/lib)
-       # MinGW TODO: Update MinGW to SDL2
-       if(NOT CMAKE_COMPILER_IS_GNUCC)
-               set(SDL_LIBRARY SDL2)
-       else()
-               set(SDL_LIBRARY SDL)
-       endif()
+       set(SDL_LIBRARY SDL2)
 endif()
 
 # Audio IO
@@ -494,14 +489,14 @@ endif()
 # used in many places so include globally, like OpenGL
 blender_include_dirs_sys("${PTHREADS_INCLUDE_DIRS}")
 
-#find signtool  
-SET(ProgramFilesX86_NAME "ProgramFiles(x86)") #env dislikes the ( ) 
+#find signtool
+set(ProgramFilesX86_NAME "ProgramFiles(x86)") #env dislikes the ( )
 find_program(SIGNTOOL_EXE signtool
-HINTS
-  "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/10/bin/x86/"
-  "$ENV{ProgramFiles}/Windows Kits/10/bin/x86/"
-  "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/8.1/bin/x86/"
-  "$ENV{ProgramFiles}/Windows Kits/8.1/bin/x86/"
-  "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/8.0/bin/x86/"
-  "$ENV{ProgramFiles}/Windows Kits/8.0/bin/x86/"
+       HINTS
+               "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/10/bin/x86/"
+               "$ENV{ProgramFiles}/Windows Kits/10/bin/x86/"
+               "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/8.1/bin/x86/"
+               "$ENV{ProgramFiles}/Windows Kits/8.1/bin/x86/"
+               "$ENV{${ProgramFilesX86_NAME}}/Windows Kits/8.0/bin/x86/"
+               "$ENV{ProgramFiles}/Windows Kits/8.0/bin/x86/"
 )
index deea844034cd9eb5fe6b89897733b367d411ff65..3ac4c4c9480f041f53280d3f599847ab6e743417 100755 (executable)
@@ -145,7 +145,7 @@ def cmake_advanced_info():
     def create_eclipse_project():
         print("CMAKE_DIR %r" % CMAKE_DIR)
         if sys.platform == "win32":
-            cmd = 'cmake "%s" -G"Eclipse CDT4 - MinGW Makefiles"' % CMAKE_DIR
+            raise Exception("Error: win32 is not supported")
         else:
             if make_exe_basename.startswith(("make", "gmake")):
                 cmd = 'cmake "%s" -G"Eclipse CDT4 - Unix Makefiles"' % CMAKE_DIR
index 310eeceadf3e0cf0302a60ef0b60a361f6942145..c1a49a756a0d00fb219fe969fcd948fd7fab20db 100644 (file)
@@ -2,13 +2,14 @@
 Modal Execution
 +++++++++++++++
 
-This operator defines a :class:`Operator.modal` function which running,
-handling events until it returns ``{'FINISHED'}`` or ``{'CANCELLED'}``.
-
-Grab, Rotate, Scale and Fly-Mode are examples of modal operators.
-They are especially useful for interactive tools,
-your operator can have its own state where keys toggle options as the operator
-runs.
+This operator defines a :class:`Operator.modal` function that will keep being
+run to handle events until it returns ``{'FINISHED'}`` or ``{'CANCELLED'}``.
+
+Modal operators run every time a new event is detected, such as a mouse click
+or key press. Conversely, when no new events are detected, the modal operator
+will not run. Modal operators are especially useful for interactive tools, an
+operator can have its own state where keys toggle options as the operator runs.
+Grab, Rotate, Scale, and Fly-Mode are examples of modal operators.
 
 :class:`Operator.invoke` is used to initialize the operator as being by
 returning ``{'RUNNING_MODAL'}``, initializing the modal loop.
index ba2e6949b811a0407b7de77efe7d3b41dec9e3e1..721374cd472397346fd4967917d3db6f257f0bf9 100644 (file)
@@ -5,23 +5,25 @@
 Python API Overview
 *******************
 
-This document is to give an understanding of how Python and Blender fit together,
-covering some of the functionality that isn't obvious from reading the API reference and example scripts.
+The purpose of this document is to explain how Python and Blender fit together,
+covering some of the functionality that may not be obvious from reading the API
+references and example scripts.
 
 
 Python in Blender
 =================
 
-Blender embeds a Python interpreter which is started with Blender and stays active.
-This interpreter runs scripts to draw the user interface and is used for some of Blender's internal tools too.
+Blender has an embedded Python interpreter which is loaded when Blender is started and stays
+active while Blender is running. This interpreter runs scripts to draw the user interface
+and is used for some of Blender’s internal tools as well.
 
-This is a typical Python environment so tutorials on how to write Python scripts
-will work running the scripts in Blender too.
-Blender provides the :mod:`bpy` module to the Python interpreter.
-This module can be imported in a script and gives access to Blender data, classes, and functions.
-Scripts that deal with Blender data will need to import this module.
+Blender's embedded interpreter provides a typical Python environment, so code from tutorials
+on how to write Python scripts can also be run with Blender’s interpreter. Blender provides its
+Python modules, such as :mod:`bpy` and :mod:`mathutils`, to the embedded interpreter so they can
+be imported into a script and give access to Blender's data, classes, and functions. Scripts that
+deal with Blender data will need to import the modules to work.
 
-Here is a simple example of moving a vertex of the object named **Cube**:
+Here is a simple example which moves a vertex attached to an object named **Cube**:
 
 .. code-block:: python
 
@@ -49,15 +51,17 @@ See the :ref:`directory layout docs <blender_manual:getting-started_installing-c
 Script Loading
 ==============
 
-This may seem obvious but it's important to note the difference
-between executing a script directly or importing it as a module.
+This may seem obvious, but it is important to note the difference between
+executing a script directly and importing a script as a module.
 
-Scripts that extend Blender - define classes that exist beyond the scripts execution,
-this makes future access to these classes (to unregister for example)
-more difficult than importing as a module where class instance is kept
-in the module and can be accessed by importing that module later on.
+Extending Blender by executing a script directly means the classes that the script
+defines remain available inside Blender after the script finishes execution.
+Using scripts this way makes future access to their classes
+(to unregister them for example) more difficult compared to importing the scripts as modules.
+When a script is imported as a module, its class instances will remain
+inside the module and can be accessed later on by importing that module again.
 
-For this reason it's preferable to only use directly execute scripts that don't extend Blender by registering classes.
+For this reason it is preferable to avoid directly executing scripts that extend Blender by registering classes.
 
 
 Here are some ways to run scripts directly in Blender.
@@ -396,8 +400,8 @@ This works just as well for PropertyGroup subclasses you define yourself.
 Dynamic Defined-Classes (Advanced)
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
-In some cases the specifier for data may not be in Blender,
-renderman shader definitions for example and it may be useful to define types and remove them on the fly.
+In some cases the specifier for data may not be in Blender, renderman shader definitions
+for example, and it may be useful to define them as types and remove them on the fly.
 
 .. code-block:: python
 
@@ -420,7 +424,7 @@ renderman shader definitions for example and it may be useful to define types an
    This is an alternative syntax for class creation in Python, better suited to constructing classes dynamically.
 
 
-Calling these operators:
+To call the operators from the previous example:
 
    >>> bpy.ops.object.operator_1()
    Hello World OBJECT_OT_operator_1
index 81b2ea4d6facd2e90190e636284894fed1cff0e6..78a101e6f980844b65e02c0dafa4157d5b3c72c8 100755 (executable)
@@ -8,11 +8,9 @@
 #include <string.h>
 #include <stdlib.h>
 
-#if !defined(__MINGW32__)
 inline int strcasecmp(const char *a, const char *b) {
   return _stricmp(a,b);
 }
-#endif
 
 inline void srandom(unsigned long input) {
   srand(input);
@@ -34,7 +32,7 @@ typedef unsigned long uintptr_t;
 
 #  if _MSC_VER < 1600
 // stdint.h is not available before VS2010
-#if defined(_WIN32) && !defined(__MINGW32__)
+#if defined(_WIN32)
 /* The __intXX are built-in types of the visual complier! So we don't
    need to include anything else here.
    This typedefs should be in sync with types from MEM_sys_types.h */
diff --git a/extern/carve/patches/mingw.patch b/extern/carve/patches/mingw.patch
deleted file mode 100644 (file)
index c237edf..0000000
+++ /dev/null
@@ -1,15 +0,0 @@
-diff -r 525472fb477a include/carve/win32.h
---- a/include/carve/win32.h    Sun Jan 15 23:07:40 2012 -0500
-+++ b/include/carve/win32.h    Wed Jan 18 00:40:10 2012 +0600
-@@ -8,9 +8,11 @@
- #include <string.h>
- #include <stdlib.h>
-+#if !defined(__MINGW32__)
- inline int strcasecmp(const char *a, const char *b) {
-   return _stricmp(a,b);
- }
-+#endif
- inline void srandom(unsigned long input) {
-   srand(input);
diff --git a/extern/carve/patches/mingw_w64.patch b/extern/carve/patches/mingw_w64.patch
deleted file mode 100644 (file)
index 26a30be..0000000
+++ /dev/null
@@ -1,13 +0,0 @@
-Index: bundle.sh
-===================================================================
---- bundle.sh  (revision 45912)
-+++ bundle.sh  (working copy)
-@@ -114,7 +114,7 @@
- if env['WITH_BF_BOOST']:
-     if env['OURPLATFORM'] not in ('win32-vc', 'win64-vc'):
-         # Boost is setting as preferred collections library in the Carve code when using MSVC compiler
--        if env['OURPLATFORM'] != 'win32-mingw':
-+        if env['OURPLATFORM'] not in ('win32-mingw', 'win64-mingw'):
-             defs.append('HAVE_BOOST_UNORDERED_COLLECTIONS')
-     defs.append('CARVE_SYSTEM_BOOST')
index b7e97d68c4c340092b775c088da428f321f2f799..c5c5fd766b8a846addc97f1884bdcc1a660ac6ff 100644 (file)
@@ -1,7 +1,6 @@
 includes.patch
 win32.patch
 mesh_iterator.patch
-mingw.patch
 gcc46.patch
 clang_is_heap_fix.patch
 strict_flags.patch
index 680bceb2421e9e290b6f0687f5fd8203d9e102bf..1a5f96505320026d493786d44c0157d98699a822 100644 (file)
@@ -8,7 +8,7 @@ diff -r e82d852e4fb0 include/carve/win32.h
 -typedef char int8_t;
 -typedef short int16_t;
 -typedef long int32_t;
-+#if defined(_WIN32) && !defined(__MINGW32__)
++#if defined(_WIN32)
 +/* The __intXX are built-in types of the visual complier! So we don't
 +   need to include anything else here.
 +   This typedefs should be in sync with types from MEM_sys_types.h */
index 2d415296dac9ca2374aa54630bd35e4fdf945280..dd446613fd0bec2137c5124f26bf943212a7fe8f 100644 (file)
@@ -19,7 +19,7 @@
 #
 # ***** END LGPL LICENSE BLOCK *****
 
-remove_extra_strict_flags()
+remove_strict_flags()
 
 if(CMAKE_COMPILER_IS_GNUCC)
        remove_cc_flag("-Wunused-macros")
index b2694a285b17c9ae3a99f75b1f1d72866f59f8ae..3018fd5b316ad5a2903c3319aca9e61f148f7637 100644 (file)
@@ -239,14 +239,15 @@ 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:
-        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');
-        engine.register_pass(scene, srl, "Denoising Albedo Variance", 3, "RGB", 'COLOR');
-        engine.register_pass(scene, srl, "Denoising Depth",           1, "Z",   'VALUE');
-        engine.register_pass(scene, srl, "Denoising Depth Variance",  1, "Z",   'VALUE');
-        engine.register_pass(scene, srl, "Denoising Shadow A",        3, "XYV", 'VECTOR');
-        engine.register_pass(scene, srl, "Denoising Shadow B",        3, "XYV", 'VECTOR');
-        engine.register_pass(scene, srl, "Denoising Image",           3, "RGB", 'COLOR');
-        engine.register_pass(scene, srl, "Denoising Image Variance",  3, "RGB", 'COLOR');
\ No newline at end of file
+    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')
+        engine.register_pass(scene, srl, "Denoising Albedo Variance", 3, "RGB", 'COLOR')
+        engine.register_pass(scene, srl, "Denoising Depth",           1, "Z",   'VALUE')
+        engine.register_pass(scene, srl, "Denoising Depth Variance",  1, "Z",   'VALUE')
+        engine.register_pass(scene, srl, "Denoising Shadow A",        3, "XYV", 'VECTOR')
+        engine.register_pass(scene, srl, "Denoising Shadow B",        3, "XYV", 'VECTOR')
+        engine.register_pass(scene, srl, "Denoising Image",           3, "RGB", 'COLOR')
+        engine.register_pass(scene, srl, "Denoising Image Variance",  3, "RGB", 'COLOR')
index 51b53a43d0b9a2152ae72c105a968e2040c4d291..7f8d28e06182ccbd38207a6ba8f734eb9481d5da 100644 (file)
@@ -695,7 +695,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)
 
@@ -1166,6 +1170,12 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
     def unregister(cls):
         del bpy.types.Scene.cycles_curves
 
+def update_render_passes(self, context):
+    scene = context.scene
+    rd = scene.render
+    rl = rd.layers.active
+    rl.update_render_passes()
+
 class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
     @classmethod
     def register(cls):
@@ -1178,27 +1188,32 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
                 name="Debug BVH Traversed Nodes",
                 description="Store Debug BVH Traversed Nodes pass",
                 default=False,
+                update=update_render_passes,
                 )
         cls.pass_debug_bvh_traversed_instances = BoolProperty(
                 name="Debug BVH Traversed Instances",
                 description="Store Debug BVH Traversed Instances pass",
                 default=False,
+                update=update_render_passes,
                 )
         cls.pass_debug_bvh_intersections = BoolProperty(
                 name="Debug BVH Intersections",
                 description="Store Debug BVH Intersections",
                 default=False,
+                update=update_render_passes,
                 )
         cls.pass_debug_ray_bounces = BoolProperty(
                 name="Debug Ray Bounces",
                 description="Store Debug Ray Bounces pass",
                 default=False,
+                update=update_render_passes,
                 )
 
         cls.use_denoising = BoolProperty(
                 name="Use Denoising",
                 description="Denoise the rendered image",
                 default=False,
+                update=update_render_passes,
                 )
         cls.denoising_diffuse_direct = BoolProperty(
                 name="Diffuse Direct",
@@ -1267,6 +1282,7 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
                 name="Store denoising passes",
                 description="Store the denoising feature passes and the noisy image",
                 default=False,
+                update=update_render_passes,
         )
 
     @classmethod
index 4ed3ccd9a2cf625eca95d677597bb0ef3e628d56..38a39e1900350a090585b904123371dfc7b6cec9 100644 (file)
@@ -531,17 +531,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):
@@ -1710,7 +1710,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)
index f13b9db7013a210297ab67f43c63ac0e5702c281..3ebe2d8cf3422839092b4a9936e3cd6cfc7407a5 100644 (file)
@@ -14,7 +14,6 @@
  * limitations under the License.
  */
 
 #include "render/mesh.h"
 #include "render/object.h"
 #include "render/scene.h"
@@ -293,7 +292,7 @@ static void create_mesh_volume_attribute(BL::Object& b_ob,
 
        if(!b_domain)
                return;
-       
+
        Attribute *attr = mesh->attributes.add(std);
        VoxelAttribute *volume_data = attr->data_voxel();
        bool is_float, is_linear;
@@ -982,7 +981,7 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
                else
                        used_shaders.push_back(scene->default_surface);
        }
-       
+
        /* test if we need to sync */
        int requested_geometry_flags = Mesh::GEOMETRY_NONE;
        if(render_layer.use_surfaces) {
@@ -1017,12 +1016,12 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
        /* ensure we only sync instanced meshes once */
        if(mesh_synced.find(mesh) != mesh_synced.end())
                return mesh;
-       
+
        mesh_synced.insert(mesh);
 
        /* create derived mesh */
        array<int> oldtriangle = mesh->triangles;
-       
+
        /* compares curve_keys rather than strands in order to handle quick hair
         * adjustments in dynamic BVH - other methods could probably do this better*/
        array<float3> oldcurve_keys = mesh->curve_keys;
@@ -1111,7 +1110,7 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
                if(memcmp(&oldcurve_radius[0], &mesh->curve_radius[0], sizeof(float)*oldcurve_radius.size()) != 0)
                        rebuild = true;
        }
-       
+
        mesh->tag_update(scene, rebuild);
 
        return mesh;
@@ -1140,7 +1139,7 @@ void BlenderSync::sync_mesh_motion(BL::Object& b_ob,
        if(scene->need_motion() == Scene::MOTION_BLUR) {
                if(!mesh->use_motion_blur)
                        return;
-               
+
                /* see if this mesh needs motion data at this time */
                vector<float> object_times = object->motion_times();
                bool found = false;
@@ -1172,7 +1171,7 @@ void BlenderSync::sync_mesh_motion(BL::Object& b_ob,
 
        if(!numverts && !numkeys)
                return;
-       
+
        /* skip objects without deforming modifiers. this is not totally reliable,
         * would need a more extensive check to see which objects are animated */
        BL::Mesh b_mesh(PointerRNA_NULL);
index c6a595775071e82b494901b9b11148d2bbba5081..2b5dd5eadea636141a6276bd9ba3082fabf2257f 100644 (file)
@@ -399,14 +399,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");
index 08ba535f282a049ffc904eca8c3ace377163aa47..3a00384458a74919fc4c51031ac6a9f0d29c759f 100644 (file)
@@ -329,6 +329,9 @@ void BlenderSync::sync_integrator()
                        integrator->ao_bounces = get_int(cscene, "ao_bounces_render");
                }
        }
+       else {
+               integrator->ao_bounces = 0;
+       }
 
        if(integrator->modified(previntegrator))
                integrator->tag_update(scene);
@@ -550,11 +553,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;
 
@@ -569,7 +577,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());
@@ -598,8 +608,6 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
                b_engine.add_pass("Debug Ray Bounces", 1, "X", b_srlay.name().c_str());
                Pass::add(PASS_RAY_BOUNCES, passes);
        }
-#else
-       (void) b_srlay;  /* Ignored. */
 #endif
 
        return passes;
index 0950285d976dfd372ef052f30763b9555874519b..4ec46424b5a4303aa0342dc58d423018e54601f6 100644 (file)
@@ -68,7 +68,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,
index abdbb6be0fd25ed7ec61a7c5217b9aead8a17663..ebbf325f95b1a0d50ad9873c067200b0d290f494 100644 (file)
@@ -299,7 +299,7 @@ static inline uint get_layer(const BL::Array<int, 20>& array)
        for(uint i = 0; i < 20; i++)
                if(array[i])
                        layer |= (1 << i);
-       
+
        return layer;
 }
 
@@ -434,7 +434,7 @@ static inline string get_string(PointerRNA& ptr, const char *name)
        string str(cstr);
        if(cstr != cstrbuf)
                MEM_freeN(cstr);
-       
+
        return str;
 }
 
@@ -451,7 +451,7 @@ static inline string blender_absolute_path(BL::BlendData& b_data,
 {
        if(path.size() >= 2 && path[0] == '/' && path[1] == '/') {
                string dirname;
-               
+
                if(b_id.library()) {
                        BL::ID b_library_id(b_id.library());
                        dirname = blender_absolute_path(b_data,
@@ -544,7 +544,7 @@ static inline BL::SmokeDomainSettings object_smoke_domain_find(BL::Object& b_ob)
                                return b_smd.domain_settings();
                }
        }
-       
+
        return BL::SmokeDomainSettings(PointerRNA_NULL);
 }
 
@@ -816,4 +816,3 @@ protected:
 CCL_NAMESPACE_END
 
 #endif /* __BLENDER_UTIL_H__ */
-
index 0603ecb3afb685c9627453edbfd6ec1923157c5b..a54bb77f9f32588fea1b711dff46a71122f26d2e 100644 (file)
@@ -68,6 +68,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;
 }
 
index 527940e8f50b185a03d38709f11d6f8a4259f95f..b3b693c630c1812b9ee74163e7dc4cbf205fc465 100644 (file)
@@ -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;
        }
 };
index e219ce3cb8e8ee0560b9183d492db85d1a6c0f03..29bb1f91a40b1fc330bc001341084c4d74fa4447 100644 (file)
@@ -176,6 +176,7 @@ public:
 
        KernelFunctions<void(*)(int, TilesInfo*, int, int, float*, float*, float*, float*, float*, int*, int, int, bool)> filter_divide_shadow_kernel;
        KernelFunctions<void(*)(int, TilesInfo*, int, int, int, int, float*, float*, int*, int, int, bool)>               filter_get_feature_kernel;
+       KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)>                                     filter_detect_outliers_kernel;
        KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)>                                     filter_combine_halves_kernel;
 
        KernelFunctions<void(*)(int, int, float*, float*, float*, int*, int, int, float, float)> filter_nlm_calc_difference_kernel;
@@ -184,9 +185,9 @@ public:
        KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int, int)>       filter_nlm_update_output_kernel;
        KernelFunctions<void(*)(float*, float*, int*, int)>                                      filter_nlm_normalize_kernel;
 
-       KernelFunctions<void(*)(float*, int, int, int, float*, int*, int*, int, int, float)>                                              filter_construct_transform_kernel;
-       KernelFunctions<void(*)(int, int, float*, float*, float*, float*, float*, int*, float*, float3*, int*, int*, int, int, int, int)> filter_nlm_construct_gramian_kernel;
-       KernelFunctions<void(*)(int, int, int, int, int, float*, int*, float*, float3*, int*, int)>                                       filter_finalize_kernel;
+       KernelFunctions<void(*)(float*, int, int, int, float*, int*, int*, int, int, float)>                              filter_construct_transform_kernel;
+       KernelFunctions<void(*)(int, int, float*, float*, float*, int*, float*, float3*, int*, int*, int, int, int, int)> filter_nlm_construct_gramian_kernel;
+       KernelFunctions<void(*)(int, int, int, int, int, float*, int*, float*, float3*, int*, int)>                       filter_finalize_kernel;
 
        KernelFunctions<void(*)(KernelGlobals *, ccl_constant KernelData*, ccl_global void*, int, ccl_global char*,
                               ccl_global uint*, int, int, int, int, int, int, int, int, ccl_global int*, int,
@@ -210,6 +211,7 @@ public:
          REGISTER_KERNEL(shader),
          REGISTER_KERNEL(filter_divide_shadow),
          REGISTER_KERNEL(filter_get_feature),
+         REGISTER_KERNEL(filter_detect_outliers),
          REGISTER_KERNEL(filter_combine_halves),
          REGISTER_KERNEL(filter_nlm_calc_difference),
          REGISTER_KERNEL(filter_nlm_blur),
@@ -463,8 +465,6 @@ public:
 
        bool denoising_reconstruct(device_ptr color_ptr,
                                   device_ptr color_variance_ptr,
-                                  device_ptr guide_ptr,
-                                  device_ptr guide_variance_ptr,
                                   device_ptr output_ptr,
                                   DenoisingTask *task)
        {
@@ -483,8 +483,8 @@ public:
                                             task->reconstruction_state.source_w - max(0, dx),
                                             task->reconstruction_state.source_h - max(0, dy)};
                        filter_nlm_calc_difference_kernel()(dx, dy,
-                                                           (float*) guide_ptr,
-                                                           (float*) guide_variance_ptr,
+                                                           (float*) color_ptr,
+                                                           (float*) color_variance_ptr,
                                                            difference,
                                                            local_rect,
                                                            task->buffer.w,
@@ -497,8 +497,6 @@ public:
                        filter_nlm_construct_gramian_kernel()(dx, dy,
                                                              blurDifference,
                                                              (float*)  task->buffer.mem.device_pointer,
-                                                             (float*)  color_ptr,
-                                                             (float*)  color_variance_ptr,
                                                              (float*)  task->storage.transform.device_pointer,
                                                              (int*)    task->storage.rank.device_pointer,
                                                              (float*)  task->storage.XtWX.device_pointer,
@@ -530,9 +528,8 @@ public:
 
        bool denoising_combine_halves(device_ptr a_ptr, device_ptr b_ptr,
                                      device_ptr mean_ptr, device_ptr variance_ptr,
-                                     int r, int4 rect, DenoisingTask *task)
+                                     int r, int4 rect, DenoisingTask * /*task*/)
        {
-               (void) task;
                for(int y = rect.y; y < rect.w; y++) {
                        for(int x = rect.x; x < rect.z; x++) {
                                filter_combine_halves_kernel()(x, y,
@@ -594,6 +591,26 @@ public:
                return true;
        }
 
+       bool denoising_detect_outliers(device_ptr image_ptr,
+                                      device_ptr variance_ptr,
+                                      device_ptr depth_ptr,
+                                      device_ptr output_ptr,
+                                      DenoisingTask *task)
+       {
+               for(int y = task->rect.y; y < task->rect.w; y++) {
+                       for(int x = task->rect.x; x < task->rect.z; x++) {
+                               filter_detect_outliers_kernel()(x, y,
+                                                               (float*) image_ptr,
+                                                               (float*) variance_ptr,
+                                                               (float*) depth_ptr,
+                                                               (float*) output_ptr,
+                                                               &task->rect.x,
+                                                               task->buffer.pass_stride);
+                       }
+               }
+               return true;
+       }
+
        void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
        {
                float *render_buffer = (float*)tile.buffer;
@@ -627,11 +644,12 @@ public:
                DenoisingTask denoising(this);
 
                denoising.functions.construct_transform = function_bind(&CPUDevice::denoising_construct_transform, this, &denoising);
-               denoising.functions.reconstruct = function_bind(&CPUDevice::denoising_reconstruct, this, _1, _2, _3, _4, _5, &denoising);
+               denoising.functions.reconstruct = function_bind(&CPUDevice::denoising_reconstruct, this, _1, _2, _3, &denoising);
                denoising.functions.divide_shadow = function_bind(&CPUDevice::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
                denoising.functions.non_local_means = function_bind(&CPUDevice::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
                denoising.functions.combine_halves = function_bind(&CPUDevice::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
                denoising.functions.get_feature = function_bind(&CPUDevice::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.detect_outliers = function_bind(&CPUDevice::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
                denoising.functions.set_tiles = function_bind(&CPUDevice::denoising_set_tiles, this, _1, &denoising);
 
                denoising.filter_area = make_int4(tile.x, tile.y, tile.w, tile.h);
index 968ee5bc487e6e664e6a1a3f0d3c033732c1c976..29fa08d94b1838a42795207d951b61dc4f591df8 100644 (file)
@@ -225,6 +225,9 @@ public:
                cuDevice = 0;
                cuContext = 0;
 
+               cuModule = 0;
+               cuFilterModule = 0;
+
                split_kernel = NULL;
 
                need_bindless_mapping = false;
@@ -487,6 +490,16 @@ public:
 
        bool load_kernels(const DeviceRequestedFeatures& requested_features)
        {
+               /* TODO(sergey): Support kernels re-load for CUDA devices.
+                *
+                * Currently re-loading kernel will invalidate memory pointers,
+                * causing problems in cuCtxSynchronize.
+                */
+               if(cuFilterModule && cuModule) {
+                       VLOG(1) << "Skipping kernel reload, not currently supported.";
+                       return true;
+               }
+
                /* check if cuda init succeeded */
                if(cuContext == 0)
                        return false;
@@ -949,7 +962,7 @@ public:
                cuda_push_context();
 
                int4 rect = task->rect;
-               int w = rect.z-rect.x;
+               int w = align_up(rect.z-rect.x, 4);
                int h = rect.w-rect.y;
                int r = task->nlm_state.r;
                int f = task->nlm_state.f;
@@ -1038,8 +1051,6 @@ public:
 
        bool denoising_reconstruct(device_ptr color_ptr,
                                   device_ptr color_variance_ptr,
-                                  device_ptr guide_ptr,
-                                  device_ptr guide_variance_ptr,
                                   device_ptr output_ptr,
                                   DenoisingTask *task)
        {
@@ -1083,8 +1094,8 @@ public:
                                             task->reconstruction_state.source_h - max(0, dy)};
 
                        void *calc_difference_args[] = {&dx, &dy,
-                                                       &guide_ptr,
-                                                       &guide_variance_ptr,
+                                                       &color_ptr,
+                                                       &color_variance_ptr,
                                                        &difference,
                                                        &local_rect,
                                                        &task->buffer.w,
@@ -1113,8 +1124,6 @@ public:
                        void *construct_gramian_args[] = {&dx, &dy,
                                                          &blurDifference,
                                                          &task->buffer.mem.device_pointer,
-                                                         &color_ptr,
-                                                         &color_variance_ptr,
                                                          &task->storage.transform.device_pointer,
                                                          &task->storage.rank.device_pointer,
                                                          &task->storage.XtWX.device_pointer,
@@ -1148,8 +1157,6 @@ public:
                                      device_ptr mean_ptr, device_ptr variance_ptr,
                                      int r, int4 rect, DenoisingTask *task)
        {
-               (void) task;
-
                if(have_error())
                        return false;
 
@@ -1179,8 +1186,6 @@ public:
                                     device_ptr sample_variance_ptr, device_ptr sv_variance_ptr,
                                     device_ptr buffer_variance_ptr, DenoisingTask *task)
        {
-               (void) task;
-
                if(have_error())
                        return false;
 
@@ -1248,16 +1253,49 @@ public:
                return !have_error();
        }
 
+       bool denoising_detect_outliers(device_ptr image_ptr,
+                                      device_ptr variance_ptr,
+                                      device_ptr depth_ptr,
+                                      device_ptr output_ptr,
+                                      DenoisingTask *task)
+       {
+               if(have_error())
+                       return false;
+
+               cuda_push_context();
+
+               CUfunction cuFilterDetectOutliers;
+               cuda_assert(cuModuleGetFunction(&cuFilterDetectOutliers, cuFilterModule, "kernel_cuda_filter_detect_outliers"));
+               cuda_assert(cuFuncSetCacheConfig(cuFilterDetectOutliers, CU_FUNC_CACHE_PREFER_L1));
+               CUDA_GET_BLOCKSIZE(cuFilterDetectOutliers,
+                                  task->rect.z-task->rect.x,
+                                  task->rect.w-task->rect.y);
+
+               void *args[] = {&image_ptr,
+                               &variance_ptr,
+                               &depth_ptr,
+                               &output_ptr,
+                               &task->rect,
+                               &task->buffer.pass_stride};
+
+               CUDA_LAUNCH_KERNEL(cuFilterDetectOutliers, args);
+               cuda_assert(cuCtxSynchronize());
+
+               cuda_pop_context();
+               return !have_error();
+       }
+
        void denoise(RenderTile &rtile, const DeviceTask &task)
        {
                DenoisingTask denoising(this);
 
                denoising.functions.construct_transform = function_bind(&CUDADevice::denoising_construct_transform, this, &denoising);
-               denoising.functions.reconstruct = function_bind(&CUDADevice::denoising_reconstruct, this, _1, _2, _3, _4, _5, &denoising);
+               denoising.functions.reconstruct = function_bind(&CUDADevice::denoising_reconstruct, this, _1, _2, _3, &denoising);
                denoising.functions.divide_shadow = function_bind(&CUDADevice::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
                denoising.functions.non_local_means = function_bind(&CUDADevice::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
                denoising.functions.combine_halves = function_bind(&CUDADevice::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
                denoising.functions.get_feature = function_bind(&CUDADevice::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.detect_outliers = function_bind(&CUDADevice::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
                denoising.functions.set_tiles = function_bind(&CUDADevice::denoising_set_tiles, this, _1, &denoising);
 
                denoising.filter_area = make_int4(rtile.x, rtile.y, rtile.w, rtile.h);
index 39c8cf301058ef0452ccc982515cd1d370aa0a55..619cc1d171ec873a1308e42f2a1ae93a8d8c82ad 100644 (file)
@@ -139,9 +139,9 @@ bool DenoisingTask::run_denoising()
                nlm_state.temporary_2_ptr = *nlm_temporary_2;
                nlm_state.temporary_3_ptr = *nlm_temporary_3;
 
-               int mean_from[]     = { 0, 1, 2, 6,  7,  8, 12 };
-               int variance_from[] = { 3, 4, 5, 9, 10, 11, 13 };
-               int pass_to[]       = { 1, 2, 3, 0,  5,  6,  7 };
+               int mean_from[]     = { 0, 1, 2, 12, 6,  7, 8 };
+               int variance_from[] = { 3, 4, 5, 13, 9, 10, 11};
+               int pass_to[]       = { 1, 2, 3, 0,  5,  6,  7};
                for(int pass = 0; pass < 7; pass++) {
                        device_sub_ptr feature_pass(device, buffer.mem, pass_to[pass]*buffer.pass_stride, buffer.pass_stride, MEM_READ_WRITE);
                        /* Get the unfiltered pass and its variance from the RenderBuffers. */
@@ -159,11 +159,25 @@ bool DenoisingTask::run_denoising()
                int mean_to[]       = { 8,  9, 10};
                int variance_to[]   = {11, 12, 13};
                int num_color_passes = 3;
+
+               device_only_memory<float> temp_color;
+               temp_color.resize(3*buffer.pass_stride);
+               device->mem_alloc("Denoising temporary color", temp_color, MEM_READ_WRITE);
+
                for(int pass = 0; pass < num_color_passes; pass++) {
-                       device_sub_ptr color_pass    (device, buffer.mem,     mean_to[pass]*buffer.pass_stride, buffer.pass_stride, MEM_READ_WRITE);
+                       device_sub_ptr color_pass(device, temp_color, pass*buffer.pass_stride, buffer.pass_stride, MEM_READ_WRITE);
                        device_sub_ptr color_var_pass(device, buffer.mem, variance_to[pass]*buffer.pass_stride, buffer.pass_stride, MEM_READ_WRITE);
                        functions.get_feature(mean_from[pass], variance_from[pass], *color_pass, *color_var_pass);
                }
+
+               {
+                       device_sub_ptr depth_pass    (device, buffer.mem,                                 0,   buffer.pass_stride, MEM_READ_WRITE);
+                       device_sub_ptr color_var_pass(device, buffer.mem, variance_to[0]*buffer.pass_stride, 3*buffer.pass_stride, MEM_READ_WRITE);
+                       device_sub_ptr output_pass   (device, buffer.mem,     mean_to[0]*buffer.pass_stride, 3*buffer.pass_stride, MEM_READ_WRITE);
+                       functions.detect_outliers(temp_color.device_pointer, *color_var_pass, *depth_pass, *output_pass);
+               }
+
+               device->mem_free(temp_color);
        }
 
        storage.w = filter_area.z;
@@ -201,7 +215,7 @@ bool DenoisingTask::run_denoising()
        {
                device_sub_ptr color_ptr    (device, buffer.mem,  8*buffer.pass_stride, 3*buffer.pass_stride, MEM_READ_WRITE);
                device_sub_ptr color_var_ptr(device, buffer.mem, 11*buffer.pass_stride, 3*buffer.pass_stride, MEM_READ_WRITE);
-               functions.reconstruct(*color_ptr, *color_var_ptr, *color_ptr, *color_var_ptr, render_buffer.ptr);
+               functions.reconstruct(*color_ptr, *color_var_ptr, render_buffer.ptr);
        }
 
        device->mem_free(storage.XtWX);
index 86d8eb643860843049b367c4eb575e242c11c6e4..def7b72f67dac23e6add695bbef3e461281dda99 100644 (file)
@@ -58,8 +58,6 @@ public:
                              )> non_local_means;
                function<bool(device_ptr color_ptr,
                              device_ptr color_variance_ptr,
-                             device_ptr guide_ptr,
-                             device_ptr guide_variance_ptr,
                              device_ptr output_ptr
                              )> reconstruct;
                function<bool()> construct_transform;
@@ -82,6 +80,11 @@ public:
                              device_ptr mean_ptr,
                              device_ptr variance_ptr
                              )> get_feature;
+               function<bool(device_ptr image_ptr,
+                             device_ptr variance_ptr,
+                             device_ptr depth_ptr,
+                             device_ptr output_ptr
+                             )> detect_outliers;
                function<bool(device_ptr*)> set_tiles;
        } functions;
 
index edd2047debcb7943c34a297d7ba8bb935f19bef7..681b8214b03fd8e0fbd8013475888eb57cc00ac4 100644 (file)
@@ -130,10 +130,22 @@ string device_opencl_capabilities(void)
                opencl_assert(func(id, what, sizeof(data), &data, NULL)); \
                result += string_printf("%s: %s\n", name, data); \
        } while(false)
+#define APPEND_STRING_EXTENSION_INFO(func, id, name, what) \
+       do { \
+               char data[1024] = "\0"; \
+               size_t length = 0; \
+               if(func(id, what, sizeof(data), &data, &length) == CL_SUCCESS) { \
+                       if(length != 0 && data[0] != '\0') { \
+                               result += string_printf("%s: %s\n", name, data); \
+                       } \
+               } \
+       } while(false)
 #define APPEND_PLATFORM_STRING_INFO(id, name, what) \
        APPEND_STRING_INFO(clGetPlatformInfo, id, "\tPlatform " name, what)
 #define APPEND_DEVICE_STRING_INFO(id, name, what) \
        APPEND_STRING_INFO(clGetDeviceInfo, id, "\t\t\tDevice " name, what)
+#define APPEND_DEVICE_STRING_EXTENSION_INFO(id, name, what) \
+       APPEND_STRING_EXTENSION_INFO(clGetDeviceInfo, id, "\t\t\tDevice " name, what)
 
        vector<cl_device_id> device_ids;
        for(cl_uint platform = 0; platform < num_platforms; ++platform) {
@@ -167,6 +179,7 @@ string device_opencl_capabilities(void)
                        result += string_printf("\t\tDevice: #%u\n", device);
 
                        APPEND_DEVICE_STRING_INFO(device_id, "Name", CL_DEVICE_NAME);
+                       APPEND_DEVICE_STRING_EXTENSION_INFO(device_id, "Board Name", CL_DEVICE_BOARD_NAME_AMD);
                        APPEND_DEVICE_STRING_INFO(device_id, "Vendor", CL_DEVICE_VENDOR);
                        APPEND_DEVICE_STRING_INFO(device_id, "OpenCL C Version", CL_DEVICE_OPENCL_C_VERSION);
                        APPEND_DEVICE_STRING_INFO(device_id, "Profile", CL_DEVICE_PROFILE);
index dddd19f179f2587b471467f2b899e8db60a3a2b6..d2b3a89fa98c329e5545d9472024c585b8b061d3 100644 (file)
@@ -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);
index 68c2ba974a5d12aefa6b980c27b47f3cc6eca480..2bac1998cb7ced6d4223df401a85256a52160f04 100644 (file)
@@ -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;
index a458ca6bf646bf7b0b1f3b5dd594e577e24873df..399fae9b42e75ec5451c3b6ad254d530fcf9dc04 100644 (file)
@@ -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.
@@ -390,8 +395,6 @@ protected:
        bool denoising_construct_transform(DenoisingTask *task);
        bool denoising_reconstruct(device_ptr color_ptr,
                                   device_ptr color_variance_ptr,
-                                  device_ptr guide_ptr,
-                                  device_ptr guide_variance_ptr,
                                   device_ptr output_ptr,
                                   DenoisingTask *task);
        bool denoising_combine_halves(device_ptr a_ptr,
@@ -411,6 +414,11 @@ protected:
                                   device_ptr mean_ptr,
                                   device_ptr variance_ptr,
                                   DenoisingTask *task);
+       bool denoising_detect_outliers(device_ptr image_ptr,
+                                      device_ptr variance_ptr,
+                                      device_ptr depth_ptr,
+                                      device_ptr output_ptr,
+                                      DenoisingTask *task);
        bool denoising_set_tiles(device_ptr *buffers,
                                 DenoisingTask *task);
 
index ae1a7b917c39b8c89da3a2c1483b9268f85f354e..e4ab979dcbf9bfce79a4bbbbdb7b28916cf7b59b 100644 (file)
@@ -216,6 +216,7 @@ bool OpenCLDeviceBase::load_kernels(const DeviceRequestedFeatures& requested_fea
        denoising_program = OpenCLProgram(this, "denoising", "filter.cl", "");
        denoising_program.add_kernel(ustring("filter_divide_shadow"));
        denoising_program.add_kernel(ustring("filter_get_feature"));
+       denoising_program.add_kernel(ustring("filter_detect_outliers"));
        denoising_program.add_kernel(ustring("filter_combine_halves"));
        denoising_program.add_kernel(ustring("filter_construct_transform"));
        denoising_program.add_kernel(ustring("filter_nlm_calc_difference"));
@@ -692,8 +693,6 @@ bool OpenCLDeviceBase::denoising_construct_transform(DenoisingTask *task)
 
 bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
                                              device_ptr color_variance_ptr,
-                                             device_ptr guide_ptr,
-                                             device_ptr guide_variance_ptr,
                                              device_ptr output_ptr,
                                              DenoisingTask *task)
 {
@@ -702,8 +701,6 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
 
        cl_mem color_mem = CL_MEM_PTR(color_ptr);
        cl_mem color_variance_mem = CL_MEM_PTR(color_variance_ptr);
-       cl_mem guide_mem = CL_MEM_PTR(guide_ptr);
-       cl_mem guide_variance_mem = CL_MEM_PTR(guide_variance_ptr);
        cl_mem output_mem = CL_MEM_PTR(output_ptr);
 
        cl_mem buffer_mem = CL_MEM_PTR(task->buffer.mem.device_pointer);
@@ -734,8 +731,8 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
 
                kernel_set_args(ckNLMCalcDifference, 0,
                                dx, dy,
-                               guide_mem,
-                               guide_variance_mem,
+                               color_mem,
+                               color_variance_mem,
                                difference,
                                local_rect,
                                task->buffer.w,
@@ -774,8 +771,6 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
                                dx, dy,
                                blurDifference,
                                buffer_mem,
-                               color_mem,
-                               color_variance_mem,
                                transform_mem,
                                rank_mem,
                                XtWX_mem,
@@ -816,8 +811,6 @@ bool OpenCLDeviceBase::denoising_combine_halves(device_ptr a_ptr,
                                                 int r, int4 rect,
                                                 DenoisingTask *task)
 {
-       (void) task;
-
        cl_mem a_mem = CL_MEM_PTR(a_ptr);
        cl_mem b_mem = CL_MEM_PTR(b_ptr);
        cl_mem mean_mem = CL_MEM_PTR(mean_ptr);
@@ -846,8 +839,6 @@ bool OpenCLDeviceBase::denoising_divide_shadow(device_ptr a_ptr,
                                                device_ptr buffer_variance_ptr,
                                                DenoisingTask *task)
 {
-       (void) task;
-
        cl_mem a_mem = CL_MEM_PTR(a_ptr);
        cl_mem b_mem = CL_MEM_PTR(b_ptr);
        cl_mem sample_variance_mem = CL_MEM_PTR(sample_variance_ptr);
@@ -910,6 +901,33 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset,
        return true;
 }
 
+bool OpenCLDeviceBase::denoising_detect_outliers(device_ptr image_ptr,
+                                                 device_ptr variance_ptr,
+                                                 device_ptr depth_ptr,
+                                                 device_ptr output_ptr,
+                                                 DenoisingTask *task)
+{
+       cl_mem image_mem = CL_MEM_PTR(image_ptr);
+       cl_mem variance_mem = CL_MEM_PTR(variance_ptr);
+       cl_mem depth_mem = CL_MEM_PTR(depth_ptr);
+       cl_mem output_mem = CL_MEM_PTR(output_ptr);
+
+       cl_kernel ckFilterDetectOutliers = denoising_program(ustring("filter_detect_outliers"));
+
+       kernel_set_args(ckFilterDetectOutliers, 0,
+                       image_mem,
+                       variance_mem,
+                       depth_mem,
+                       output_mem,
+                       task->rect,
+                       task->buffer.pass_stride);
+       enqueue_kernel(ckFilterDetectOutliers,
+                      task->rect.z-task->rect.x,
+                      task->rect.w-task->rect.y);
+
+       return true;
+}
+
 bool OpenCLDeviceBase::denoising_set_tiles(device_ptr *buffers,
                                            DenoisingTask *task)
 {
@@ -937,11 +955,12 @@ void OpenCLDeviceBase::denoise(RenderTile &rtile, const DeviceTask &task)
 
        denoising.functions.set_tiles = function_bind(&OpenCLDeviceBase::denoising_set_tiles, this, _1, &denoising);
        denoising.functions.construct_transform = function_bind(&OpenCLDeviceBase::denoising_construct_transform, this, &denoising);
-       denoising.functions.reconstruct = function_bind(&OpenCLDeviceBase::denoising_reconstruct, this, _1, _2, _3, _4, _5, &denoising);
+       denoising.functions.reconstruct = function_bind(&OpenCLDeviceBase::denoising_reconstruct, this, _1, _2, _3, &denoising);
        denoising.functions.divide_shadow = function_bind(&OpenCLDeviceBase::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
        denoising.functions.non_local_means = function_bind(&OpenCLDeviceBase::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
        denoising.functions.combine_halves = function_bind(&OpenCLDeviceBase::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
        denoising.functions.get_feature = function_bind(&OpenCLDeviceBase::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+       denoising.functions.detect_outliers = function_bind(&OpenCLDeviceBase::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
 
        denoising.filter_area = make_int4(rtile.x, rtile.y, rtile.w, rtile.h);
        denoising.render_buffer.samples = rtile.sample;
index 76dcbd6fc9adeb7143b428a75fba481863e3e317..08b632ee9d3670669801b818dabd22424f7a7e44 100644 (file)
@@ -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;
        }
 
index 642c1bfa11c9bb72c3dec99e287dd9c1a5e34d4e..8ba2a8e26da18873f3f8f5da508cee0b73e74a4b 100644 (file)
@@ -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) {
@@ -1063,7 +1086,7 @@ string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
                           CL_DEVICE_BOARD_NAME_AMD,
                           sizeof(board_name),
                           &board_name,
-                                          &length) == CL_SUCCESS)
+                          &length) == CL_SUCCESS)
        {
                if(length != 0 && board_name[0] != '\0') {
                        return board_name;
@@ -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;
index bef869f34b4da20763c837c76679b2874050931c..23e9bd311c45eb3704dc00ff937defc833c85052 100644 (file)
@@ -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)
index 30644946840d0171c552a26731a4f56b1a972ece..b07b515c405f302953c5fb83d6d6b1e24ef5ae1f 100644 (file)
@@ -91,18 +91,15 @@ ccl_device_forceinline float3 mf_sample_vndf(const float3 wi, const float2 alpha
        return normalize(make_float3(-slope_x, -slope_y, 1.0f));
 }
 
-/* === Phase functions: Glossy, Diffuse and Glass === */
+/* === Phase functions: Glossy and Glass === */
 
-/* Phase function for reflective materials, either without a fresnel term (for compatibility) or with the conductive fresnel term. */
-ccl_device_forceinline float3 mf_sample_phase_glossy(const float3 wi, float3 *n, float3 *k, float3 *weight, const float3 wm)
+/* Phase function for reflective materials. */
+ccl_device_forceinline float3 mf_sample_phase_glossy(const float3 wi, float3 *weight, const float3 wm)
 {
-       if(n && k)
-               *weight *= fresnel_conductor(dot(wi, wm), *n, *k);
-
        return -wi + 2.0f * wm * dot(wi, wm);
 }
 
-ccl_device_forceinline float3 mf_eval_phase_glossy(const float3 w, const float lambda, const float3 wo, const float2 alpha, float3 *n, float3 *k)
+ccl_device_forceinline float3 mf_eval_phase_glossy(const float3 w, const float lambda, const float3 wo, const float2 alpha)
 {
        if(w.z > 0.9999f)
                return make_float3(0.0f, 0.0f, 0.0f);
@@ -123,30 +120,9 @@ ccl_device_forceinline float3 mf_eval_phase_glossy(const float3 w, const float l
        else
                phase *= D_ggx_aniso(wh, alpha);
 
-       if(n && k) {
-               /* Apply conductive fresnel term. */
-               return phase * fresnel_conductor(dotW_WH, *n, *k);
-       }
-
        return make_float3(phase, phase, phase);
 }
 
-/* Phase function for rough lambertian diffuse surfaces. */
-ccl_device_forceinline float3 mf_sample_phase_diffuse(const float3 wm, const float randu, const float randv)
-{
-       float3 tm, bm;
-       make_orthonormals(wm, &tm, &bm);
-
-       float2 disk = concentric_sample_disk(randu, randv);
-       return disk.x*tm + disk.y*bm + safe_sqrtf(1.0f - disk.x*disk.x - disk.y*disk.y)*wm;
-}
-
-ccl_device_forceinline float3 mf_eval_phase_diffuse(const float3 w, const float3 wm)
-{
-       const float v = max(0.0f, dot(w, wm)) * M_1_PI_F;
-       return make_float3(v, v, v);
-}
-
 /* Phase function for dielectric transmissive materials, including both reflection and refraction according to the dielectric fresnel term. */
 ccl_device_forceinline float3 mf_sample_phase_glass(const float3 wi, const float eta, const float3 wm, const float randV, bool *outside)
 {
@@ -282,11 +258,6 @@ ccl_device_forceinline float mf_ggx_aniso_pdf(const float3 wi, const float3 wo,
        return 0.25f * D_ggx_aniso(normalize(wi+wo), alpha) / ((1.0f + mf_lambda(wi, alpha)) * wi.z) + (1.0f - mf_ggx_albedo(sqrtf(alpha.x*alpha.y))) * wo.z;
 }
 
-ccl_device_forceinline float mf_diffuse_pdf(const float3 wo)
-{
-       return M_1_PI_F * wo.z;
-}
-
 ccl_device_forceinline float mf_glass_pdf(const float3 wi, const float3 wo, const float alpha, const float eta)
 {
        float3 wh;
@@ -315,13 +286,6 @@ ccl_device_forceinline float mf_glass_pdf(const float3 wi, const float3 wo, cons
 #define MF_MULTI_GLASS
 #include "kernel/closure/bsdf_microfacet_multi_impl.h"
 
-/* The diffuse phase function is not implemented as a node yet. */
-#if 0
-#define MF_PHASE_FUNCTION diffuse
-#define MF_MULTI_DIFFUSE
-#include "kernel/closure/bsdf_microfacet_multi_impl.h"
-#endif
-
 #define MF_PHASE_FUNCTION glossy
 #define MF_MULTI_GLOSSY
 #include "kernel/closure/bsdf_microfacet_multi_impl.h"
@@ -428,7 +392,7 @@ ccl_device float3 bsdf_microfacet_multi_ggx_eval_reflect(const ShaderClosure *sc
                *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
        else
                *pdf = mf_ggx_pdf(localI, localO, bsdf->alpha_x);
-       return mf_eval_glossy(localI, localO, true, bsdf->extra->color, bsdf->alpha_x, bsdf->alpha_y, lcg_state, NULL, NULL, bsdf->ior, use_fresnel, bsdf->extra->cspec0);
+       return mf_eval_glossy(localI, localO, true, bsdf->extra->color, bsdf->alpha_x, bsdf->alpha_y, lcg_state, bsdf->ior, use_fresnel, bsdf->extra->cspec0);
 }
 
 ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals *kg, const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf, ccl_addr_space uint *lcg_state)
@@ -442,6 +406,10 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals *kg, const ShaderC
                *omega_in = 2*dot(Z, I)*Z - I;
                *pdf = 1e6f;
                *eval = make_float3(1e6f, 1e6f, 1e6f);
+#ifdef __RAY_DIFFERENTIALS__
+               *domega_in_dx = (2 * dot(Z, dIdx)) * Z - dIdx;
+               *domega_in_dy = (2 * dot(Z, dIdy)) * Z - dIdy;
+#endif
                return LABEL_REFLECT|LABEL_SINGULAR;
        }
 
@@ -456,7 +424,7 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals *kg, const ShaderC
        float3 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
        float3 localO;
 
-       *eval = mf_sample_glossy(localI, &localO, bsdf->extra->color, bsdf->alpha_x, bsdf->alpha_y, lcg_state, NULL, NULL, bsdf->ior, use_fresnel, bsdf->extra->cspec0);
+       *eval = mf_sample_glossy(localI, &localO, bsdf->extra->color, bsdf->alpha_x, bsdf->alpha_y, lcg_state, bsdf->ior, use_fresnel, bsdf->extra->cspec0);
        if(is_aniso)
                *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
        else
index 16d900088cbf504aceee4b0303e7b5e24a8c298d..2eb2457c9e5634042e26cb659fab795539415c05 100644 (file)
  * the balance heuristic isn't necessarily optimal anymore.
  */
 ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
-        float3 wi,
-        float3 wo,
-        const bool wo_outside,
-        const float3 color,
-        const float alpha_x,
-        const float alpha_y,
-         ccl_addr_space uint *lcg_state
-#ifdef MF_MULTI_GLASS
-        , const float eta
-        , bool use_fresnel
-        , const float3 cspec0
-#elif defined(MF_MULTI_GLOSSY)
-         , float3 *n, float3 *k
-         , const float eta
-         , bool use_fresnel
-         , const float3 cspec0
-#endif
-)
+       float3 wi,
+       float3 wo,
+       const bool wo_outside,
+       const float3 color,
+       const float alpha_x,
+       const float alpha_y,
+       ccl_addr_space uint *lcg_state,
+       const float eta,
+       bool use_fresnel,
+       const float3 cspec0)
 {
        /* Evaluating for a shallower incoming direction produces less noise, and the properties of the BSDF guarantee reciprocity. */
        bool swapped = false;
@@ -77,44 +69,29 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
        /* Analytically compute single scattering for lower noise. */
        float3 eval;
        float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
+       const float3 wh = normalize(wi+wo);
 #ifdef MF_MULTI_GLASS
        eval = mf_eval_phase_glass(-wi, lambda_r, wo, wo_outside, alpha, eta);
        if(wo_outside)
                eval *= -lambda_r / (shadowing_lambda - lambda_r);
        else
                eval *= -lambda_r * beta(-lambda_r, shadowing_lambda+1.0f);
-
-       float F0 = fresnel_dielectric_cos(1.0f, eta);
-       if(use_fresnel) {
-               throughput = interpolate_fresnel_color(wi, normalize(wi + wo), eta, F0, cspec0);
-
-               eval *= throughput;
-       }
-#elif defined(MF_MULTI_DIFFUSE)
-       /* Diffuse has no special closed form for the single scattering bounce */
-       eval = make_float3(0.0f, 0.0f, 0.0f);
 #else /* MF_MULTI_GLOSSY */
-       const float3 wh = normalize(wi+wo);
        const float G2 = 1.0f / (1.0f - (lambda_r + 1.0f) + shadowing_lambda);
        float val = G2 * 0.25f / wi.z;
        if(alpha.x == alpha.y)
                val *= D_ggx(wh, alpha.x);
        else
                val *= D_ggx_aniso(wh, alpha);
-       if(n && k) {
-               eval = fresnel_conductor(dot(wh, wi), *n, *k) * val;
-       }
-       else {
-               eval = make_float3(val, val, val);
-       }
+       eval = make_float3(val, val, val);
+#endif
 
        float F0 = fresnel_dielectric_cos(1.0f, eta);
        if(use_fresnel) {
                throughput = interpolate_fresnel_color(wi, wh, eta, F0, cspec0);
 
-               eval = throughput * val;
+               eval *= throughput;
        }
-#endif
 
        float3 wr = -wi;
        float hr = 1.0f;
@@ -129,13 +106,6 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
                float3 wm = mf_sample_vndf(-wr, alpha, make_float2(lcg_step_float_addrspace(lcg_state),
                                                                   lcg_step_float_addrspace(lcg_state)));
 
-#ifdef MF_MULTI_DIFFUSE
-               if(order == 0) {
-                       /* Compute single-scattering for diffuse. */
-                       const float G2_G1 = -lambda_r / (shadowing_lambda - lambda_r);
-                       eval += throughput * G2_G1 * mf_eval_phase_diffuse(wo, wm);
-               }
-#endif
 #ifdef MF_MULTI_GLASS
                if(order == 0 && use_fresnel) {
                        /* Evaluate amount of scattering towards wo on this microfacet. */
@@ -156,10 +126,8 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
                                phase = mf_eval_phase_glass(wr, lambda_r,  wo,  wo_outside, alpha, eta);
                        else
                                phase = mf_eval_phase_glass(wr, lambda_r, -wo, !wo_outside, alpha, 1.0f/eta);
-#elif defined(MF_MULTI_DIFFUSE)
-                       phase = mf_eval_phase_diffuse(wo, wm);
 #else /* MF_MULTI_GLOSSY */
-                       phase = mf_eval_phase_glossy(wr, lambda_r, wo, alpha, n, k) * throughput;
+                       phase = mf_eval_phase_glossy(wr, lambda_r, wo, alpha) * throughput;
 #endif
                        eval += throughput * phase * mf_G1(wo_outside? wo: -wo, mf_C1((outside == wo_outside)? hr: -hr), shadowing_lambda);
                }
@@ -181,25 +149,17 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
                        else if(use_fresnel && order > 0) {
                                throughput *= interpolate_fresnel_color(wi_prev, wm, eta, F0, cspec0);
                        }
-#elif defined(MF_MULTI_DIFFUSE)
-                       wr = mf_sample_phase_diffuse(wm,
-                                                    lcg_step_float_addrspace(lcg_state),
-                                                    lcg_step_float_addrspace(lcg_state));
 #else /* MF_MULTI_GLOSSY */
                        if(use_fresnel && order > 0) {
                                throughput *= interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
                        }
-                       wr = mf_sample_phase_glossy(-wr, n, k, &throughput, wm);
+                       wr = mf_sample_phase_glossy(-wr, &throughput, wm);
 #endif
 
                        lambda_r = mf_lambda(wr, alpha);
 
-#if defined(MF_MULTI_GLOSSY) || defined(MF_MULTI_GLASS)
                        if(!use_fresnel)
                                throughput *= color;
-#else
-                       throughput *= color;
-#endif
 
                        C1_r = mf_C1(hr);
                        G1_r = mf_G1(wr, C1_r, lambda_r);
@@ -215,18 +175,16 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_eval)(
  * escaped the surface in wo. The function returns the throughput between wi and wo.
  * Without reflection losses due to coloring or fresnel absorption in conductors, the sampling is optimal.
  */
-ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3 *wo, const float3 color, const float alpha_x, const float alpha_y, ccl_addr_space uint *lcg_state
-#ifdef MF_MULTI_GLASS
-       , const float eta
-       , bool use_fresnel
-       , const float3 cspec0
-#elif defined(MF_MULTI_GLOSSY)
-       , float3 *n, float3 *k
-       , const float eta
-       , bool use_fresnel
-       , const float3 cspec0
-#endif
-)
+ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(
+       float3 wi,
+       float3 *wo,
+       const float3 color,
+       const float alpha_x,
+       const float alpha_y,
+       ccl_addr_space uint *lcg_state,
+       const float eta,
+       bool use_fresnel,
+       const float3 cspec0)
 {
        const float2 alpha = make_float2(alpha_x, alpha_y);
 
@@ -237,17 +195,11 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3
        float C1_r = 1.0f;
        float G1_r = 0.0f;
        bool outside = true;
-#ifdef MF_MULTI_GLASS
-       float F0 = fresnel_dielectric_cos(1.0f, eta);
-       if(use_fresnel) {
-               throughput = interpolate_fresnel_color(wi, normalize(wi + wr), eta, F0, cspec0);
-       }
-#elif defined(MF_MULTI_GLOSSY)
+
        float F0 = fresnel_dielectric_cos(1.0f, eta);
        if(use_fresnel) {
                throughput = interpolate_fresnel_color(wi, normalize(wi + wr), eta, F0, cspec0);
        }
-#endif
 
        int order;
        for(order = 0; order < 10; order++) {
@@ -262,13 +214,8 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3
                                                                   lcg_step_float_addrspace(lcg_state)));
 
                /* First-bounce color is already accounted for in mix weight. */
-#if defined(MF_MULTI_GLASS) || defined(MF_MULTI_GLOSSY)
                if(!use_fresnel && order > 0)
                        throughput *= color;
-#else
-               if(order > 0)
-                       throughput *= color;
-#endif
 
                /* Bounce from the microfacet. */
 #ifdef MF_MULTI_GLASS
@@ -294,10 +241,6 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3
                                        throughput *= t_color;
                        }
                }
-#elif defined(MF_MULTI_DIFFUSE)
-               wr = mf_sample_phase_diffuse(wm,
-                                            lcg_step_float_addrspace(lcg_state),
-                                            lcg_step_float_addrspace(lcg_state));
 #else /* MF_MULTI_GLOSSY */
                if(use_fresnel) {
                        float3 t_color = interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
@@ -307,7 +250,7 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3
                        else
                                throughput *= t_color;
                }
-               wr = mf_sample_phase_glossy(-wr, n, k, &throughput, wm);
+               wr = mf_sample_phase_glossy(-wr, &throughput, wm);
 #endif
 
                /* Update random walk parameters. */
@@ -319,6 +262,5 @@ ccl_device_forceinline float3 MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi, float3
 }
 
 #undef MF_MULTI_GLASS
-#undef MF_MULTI_DIFFUSE
 #undef MF_MULTI_GLOSSY
 #undef MF_PHASE_FUNCTION
index f5a40d4999743c9e4bd54c7d4fe4649dde08f225..6226ed2c2eff21b38c2ffaebb650dd5ae7ac1c36 100644 (file)
@@ -16,7 +16,7 @@
 
  CCL_NAMESPACE_BEGIN
 
-#define ccl_get_feature(buffer, pass) buffer[(pass)*pass_stride]
+#define ccl_get_feature(buffer, pass) (buffer)[(pass)*pass_stride]
 
 /* Loop over the pixels in the range [low.x, high.x) x [low.y, high.y).
  * pixel_buffer always points to the current pixel in the first pass. */
                                  pixel_buffer += buffer_w - (high.x - low.x); \
                              }
 
-ccl_device_inline void filter_get_features(int2 pixel, ccl_global float ccl_restrict_ptr buffer, float *features, float ccl_restrict_ptr mean, int pass_stride)
+ccl_device_inline void filter_get_features(int2 pixel,
+                                           const ccl_global float *ccl_restrict buffer,
+                                           float *features,
+                                           const float *ccl_restrict mean,
+                                           int pass_stride)
 {
        features[0] = pixel.x;
        features[1] = pixel.y;
-       features[2] = ccl_get_feature(buffer, 0);
+       features[2] = fabsf(ccl_get_feature(buffer, 0));
        features[3] = ccl_get_feature(buffer, 1);
        features[4] = ccl_get_feature(buffer, 2);
        features[5] = ccl_get_feature(buffer, 3);
@@ -46,11 +50,15 @@ ccl_device_inline void filter_get_features(int2 pixel, ccl_global float ccl_rest
        }
 }
 
-ccl_device_inline void filter_get_feature_scales(int2 pixel, ccl_global float ccl_restrict_ptr buffer, float *scales, float ccl_restrict_ptr mean, int pass_stride)
+ccl_device_inline void filter_get_feature_scales(int2 pixel,
+                                                 const ccl_global float *ccl_restrict buffer,
+                                                 float *scales,
+                                                 const float *ccl_restrict mean,
+                                                 int pass_stride)
 {
        scales[0] = fabsf(pixel.x - mean[0]);
        scales[1] = fabsf(pixel.y - mean[1]);
-       scales[2] = fabsf(ccl_get_feature(buffer, 0) - mean[2]);
+       scales[2] = fabsf(fabsf(ccl_get_feature(buffer, 0)) - mean[2]);
        scales[3] = len_squared(make_float3(ccl_get_feature(buffer, 1) - mean[3],
                                            ccl_get_feature(buffer, 2) - mean[4],
                                            ccl_get_feature(buffer, 3) - mean[5]));
@@ -70,19 +78,15 @@ ccl_device_inline void filter_calculate_scale(float *scale)
        scale[3] = scale[4] = scale[5] = 1.0f/max(sqrtf(scale[3]), 0.01f);
 }
 
-ccl_device_inline float3 filter_get_pixel_color(ccl_global float ccl_restrict_ptr buffer, int pass_stride)
+ccl_device_inline float3 filter_get_color(const ccl_global float *ccl_restrict buffer,
+                                          int pass_stride)
 {
-       return make_float3(ccl_get_feature(buffer, 0), ccl_get_feature(buffer, 1), ccl_get_feature(buffer, 2));
-}
-
-ccl_device_inline float filter_get_pixel_variance(ccl_global float ccl_restrict_ptr buffer, int pass_stride)
-{
-       return average(make_float3(ccl_get_feature(buffer, 0), ccl_get_feature(buffer, 1), ccl_get_feature(buffer, 2)));
+       return make_float3(ccl_get_feature(buffer, 8), ccl_get_feature(buffer, 9), ccl_get_feature(buffer, 10));
 }
 
 ccl_device_inline void design_row_add(float *design_row,
                                       int rank,
-                                      ccl_global float ccl_restrict_ptr transform,
+                                      const ccl_global float *ccl_restrict transform,
                                       int stride,
                                       int row,
                                       float feature)
@@ -94,20 +98,20 @@ ccl_device_inline void design_row_add(float *design_row,
 
 /* Fill the design row. */
 ccl_device_inline void filter_get_design_row_transform(int2 p_pixel,
-                                                       ccl_global float ccl_restrict_ptr p_buffer,
+                                                       const ccl_global float *ccl_restrict p_buffer,
                                                        int2 q_pixel,
-                                                       ccl_global float ccl_restrict_ptr q_buffer,
+                                                       const ccl_global float *ccl_restrict q_buffer,
                                                        int pass_stride,
                                                        int rank,
                                                        float *design_row,
-                                                       ccl_global float ccl_restrict_ptr transform,
+                                                       const ccl_global float *ccl_restrict transform,
                                                        int stride)
 {
        design_row[0] = 1.0f;
        math_vector_zero(design_row+1, rank);
        design_row_add(design_row, rank, transform, stride, 0, q_pixel.x - p_pixel.x);
        design_row_add(design_row, rank, transform, stride, 1, q_pixel.y - p_pixel.y);
-       design_row_add(design_row, rank, transform, stride, 2, ccl_get_feature(q_buffer, 0) - ccl_get_feature(p_buffer, 0));
+       design_row_add(design_row, rank, transform, stride, 2, fabsf(ccl_get_feature(q_buffer, 0)) - fabsf(ccl_get_feature(p_buffer, 0)));
        design_row_add(design_row, rank, transform, stride, 3, ccl_get_feature(q_buffer, 1) - ccl_get_feature(p_buffer, 1));
        design_row_add(design_row, rank, transform, stride, 4, ccl_get_feature(q_buffer, 2) - ccl_get_feature(p_buffer, 2));
        design_row_add(design_row, rank, transform, stride, 5, ccl_get_feature(q_buffer, 3) - ccl_get_feature(p_buffer, 3));
index 303c8f482e3ad64e2b077ced9994110ea7e87b22..3185330994c69557718219ed1646c77775a78150 100644 (file)
@@ -33,11 +33,16 @@ CCL_NAMESPACE_BEGIN
                                      pixel_buffer += buffer_w - (pixel.x - low.x); \
                                  }
 
-ccl_device_inline void filter_get_features_sse(__m128 x, __m128 y, __m128 active_pixels, float ccl_restrict_ptr buffer, __m128 *features, __m128 ccl_restrict_ptr mean, int pass_stride)
+ccl_device_inline void filter_get_features_sse(__m128 x, __m128 y,
+                                               __m128 active_pixels,
+                                               const float *ccl_restrict buffer,
+                                               __m128 *features,
+                                               const __m128 *ccl_restrict mean,
+                                               int pass_stride)
 {
        features[0] = x;
        features[1] = y;
-       features[2] = ccl_get_feature_sse(0);
+       features[2] = _mm_fabs_ps(ccl_get_feature_sse(0));
        features[3] = ccl_get_feature_sse(1);
        features[4] = ccl_get_feature_sse(2);
        features[5] = ccl_get_feature_sse(3);
@@ -53,12 +58,17 @@ ccl_device_inline void filter_get_features_sse(__m128 x, __m128 y, __m128 active
                features[i] = _mm_mask_ps(features[i], active_pixels);
 }
 
-ccl_device_inline void filter_get_feature_scales_sse(__m128 x, __m128 y, __m128 active_pixels, float ccl_restrict_ptr buffer, __m128 *scales, __m128 ccl_restrict_ptr mean, int pass_stride)
+ccl_device_inline void filter_get_feature_scales_sse(__m128 x, __m128 y,
+                                                     __m128 active_pixels,
+                                                     const float *ccl_restrict buffer,
+                                                     __m128 *scales,
+                                                     const __m128 *ccl_restrict mean,
+                                                     int pass_stride)
 {
        scales[0] = _mm_mask_ps(_mm_fabs_ps(_mm_sub_ps(x, mean[0])), active_pixels);
        scales[1] = _mm_mask_ps(_mm_fabs_ps(_mm_sub_ps(y, mean[1])), active_pixels);
 
-       scales[2] = _mm_mask_ps(_mm_fabs_ps(_mm_sub_ps(ccl_get_feature_sse(0), mean[2])), active_pixels);
+       scales[2] = _mm_mask_ps(_mm_fabs_ps(_mm_sub_ps(_mm_fabs_ps(ccl_get_feature_sse(0)), mean[2])), active_pixels);
 
        __m128 diff, scale;
        diff = _mm_sub_ps(ccl_get_feature_sse(1), mean[3]);
index 1a314b100be07f3b83bbf9cb744ef3d5016d1412..3e752bce68f01afd5bae510ad0a0b707aaee517c 100644 (file)
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy, float ccl_restrict_ptr weightImage, float ccl_restrict_ptr varianceImage, float *differenceImage, int4 rect, int w, int channel_offset, float a, float k_2)
+ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy,
+                                                         const float *ccl_restrict weight_image,
+                                                         const float *ccl_restrict variance_image,
+                                                         float *difference_image,
+                                                         int4 rect,
+                                                         int w,
+                                                         int channel_offset,
+                                                         float a,
+                                                         float k_2)
 {
        for(int y = rect.y; y < rect.w; y++) {
                for(int x = rect.x; x < rect.z; x++) {
                        float diff = 0.0f;
                        int numChannels = channel_offset? 3 : 1;
                        for(int c = 0; c < numChannels; c++) {
-                               float cdiff = weightImage[c*channel_offset + y*w+x] - weightImage[c*channel_offset + (y+dy)*w+(x+dx)];
-                               float pvar = varianceImage[c*channel_offset + y*w+x];
-                               float qvar = varianceImage[c*channel_offset + (y+dy)*w+(x+dx)];
+                               float cdiff = weight_image[c*channel_offset + y*w+x] - weight_image[c*channel_offset + (y+dy)*w+(x+dx)];
+                               float pvar = variance_image[c*channel_offset + y*w+x];
+                               float qvar = variance_image[c*channel_offset + (y+dy)*w+(x+dx)];
                                diff += (cdiff*cdiff - a*(pvar + min(pvar, qvar))) / (1e-8f + k_2*(pvar+qvar));
                        }
                        if(numChannels > 1) {
                                diff *= 1.0f/numChannels;
                        }
-                       differenceImage[y*w+x] = diff;
+                       difference_image[y*w+x] = diff;
                }
        }
 }
 
-ccl_device_inline void kernel_filter_nlm_blur(float ccl_restrict_ptr differenceImage, float *outImage, int4 rect, int w, int f)
+ccl_device_inline void kernel_filter_nlm_blur(const float *ccl_restrict difference_image,
+                                              float *out_image,
+                                              int4 rect,
+                                              int w,
+                                              int f)
 {
 #ifdef __KERNEL_SSE3__
        int aligned_lowx = (rect.x & ~(3));
@@ -46,30 +58,34 @@ ccl_device_inline void kernel_filter_nlm_blur(float ccl_restrict_ptr differenceI
                const int low = max(rect.y, y-f);
                const int high = min(rect.w, y+f+1);
                for(int x = rect.x; x < rect.z; x++) {
-                       outImage[y*w+x] = 0.0f;
+                       out_image[y*w+x] = 0.0f;
                }
                for(int y1 = low; y1 < high; y1++) {
 #ifdef __KERNEL_SSE3__
                        for(int x = aligned_lowx; x < aligned_highx; x+=4) {
-                               _mm_store_ps(outImage + y*w+x, _mm_add_ps(_mm_load_ps(outImage + y*w+x), _mm_load_ps(differenceImage + y1*w+x)));
+                               _mm_store_ps(out_image + y*w+x, _mm_add_ps(_mm_load_ps(out_image + y*w+x), _mm_load_ps(difference_image + y1*w+x)));
                        }
 #else
                        for(int x = rect.x; x < rect.z; x++) {
-                               outImage[y*w+x] += differenceImage[y1*w+x];
+                               out_image[y*w+x] += difference_image[y1*w+x];
                        }
 #endif
                }
                for(int x = rect.x; x < rect.z; x++) {
-                       outImage[y*w+x] *= 1.0f/(high - low);
+                       out_image[y*w+x] *= 1.0f/(high - low);
                }
        }
 }
 
-ccl_device_inline void kernel_filter_nlm_calc_weight(float ccl_restrict_ptr differenceImage, float *outImage, int4 rect, int w, int f)
+ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict difference_image,
+                                                     float *out_image,
+                                                     int4 rect,
+                                                     int w,
+                                                     int f)
 {
        for(int y = rect.y; y < rect.w; y++) {
                for(int x = rect.x; x < rect.z; x++) {
-                       outImage[y*w+x] = 0.0f;
+                       out_image[y*w+x] = 0.0f;
                }
        }
        for(int dx = -f; dx <= f; dx++) {
@@ -77,7 +93,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(float ccl_restrict_ptr diff
                int neg_dx = min(0, dx);
                for(int y = rect.y; y < rect.w; y++) {
                        for(int x = rect.x-neg_dx; x < rect.z-pos_dx; x++) {
-                               outImage[y*w+x] += differenceImage[y*w+dx+x];
+                               out_image[y*w+x] += difference_image[y*w+dx+x];
                        }
                }
        }
@@ -85,12 +101,19 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(float ccl_restrict_ptr diff
                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);
-                       outImage[y*w+x] = expf(-max(outImage[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));
                }
        }
 }
 
-ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy, float ccl_restrict_ptr differenceImage, float ccl_restrict_ptr image, float *outImage, float *accumImage, int4 rect, int w, int f)
+ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy,
+                                                       const float *ccl_restrict difference_image,
+                                                       const float *ccl_restrict image,
+                                                       float *out_image,
+                                                       float *accum_image,
+                                                       int4 rect,
+                                                       int w,
+                                                       int f)
 {
        for(int y = rect.y; y < rect.w; y++) {
                for(int x = rect.x; x < rect.z; x++) {
@@ -98,20 +121,18 @@ ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy, float ccl
                        const int high = min(rect.z, x+f+1);
                        float sum = 0.0f;
                        for(int x1 = low; x1 < high; x1++) {
-                               sum += differenceImage[y*w+x1];
+                               sum += difference_image[y*w+x1];
                        }
                        float weight = sum * (1.0f/(high - low));
-                       accumImage[y*w+x] += weight;
-                       outImage[y*w+x] += weight*image[(y+dy)*w+(x+dx)];
+                       accum_image[y*w+x] += weight;
+                       out_image[y*w+x] += weight*image[(y+dy)*w+(x+dx)];
                }
        }
 }
 
 ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
-                                                           float ccl_restrict_ptr differenceImage,
-                                                           float ccl_restrict_ptr buffer,
-                                                           float *color_pass,
-                                                           float *variance_pass,
+                                                           const float *ccl_restrict difference_image,
+                                                           const float *ccl_restrict buffer,
                                                            float *transform,
                                                            int *rank,
                                                            float *XtWX,
@@ -130,7 +151,7 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
                        const int high = min(rect.z, x+f+1);
                        float sum = 0.0f;
                        for(int x1 = low; x1 < high; x1++) {
-                               sum += differenceImage[y*w+x1];
+                               sum += difference_image[y*w+x1];
                        }
                        float weight = sum * (1.0f/(high - low));
 
@@ -144,18 +165,20 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
                                                        dx, dy, w, h,
                                                        pass_stride,
                                                        buffer,
-                                                       color_pass, variance_pass,
                                                        l_transform, l_rank,
                                                        weight, l_XtWX, l_XtWY, 0);
                }
        }
 }
 
-ccl_device_inline void kernel_filter_nlm_normalize(float *outImage, float ccl_restrict_ptr accumImage, int4 rect, int w)
+ccl_device_inline void kernel_filter_nlm_normalize(float *out_image,
+                                                   const float *ccl_restrict accum_image,
+                                                   int4 rect,
+                                                   int w)
 {
        for(int y = rect.y; y < rect.w; y++) {
                for(int x = rect.x; x < rect.z; x++) {
-                       outImage[y*w+x] /= accumImage[y*w+x];
+                       out_image[y*w+x] /= accum_image[y*w+x];
                }
        }
 }
index b5ba7cf51a5b22c45a0243b79895a98eacf2966f..2c5ac8070513f231ea50c26cfea3c6650c3de9cf 100644 (file)
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 ccl_device_inline void kernel_filter_nlm_calc_difference(int x, int y,
                                                          int dx, int dy,
-                                                         ccl_global float ccl_restrict_ptr weightImage,
-                                                         ccl_global float ccl_restrict_ptr varianceImage,
-                                                         ccl_global float *differenceImage,
+                                                         const ccl_global float *ccl_restrict weight_image,
+                                                         const ccl_global float *ccl_restrict variance_image,
+                                                         ccl_global float *difference_image,
                                                          int4 rect, int w,
                                                          int channel_offset,
                                                          float a, float k_2)
@@ -28,78 +28,76 @@ ccl_device_inline void kernel_filter_nlm_calc_difference(int x, int y,
        float diff = 0.0f;
        int numChannels = channel_offset? 3 : 1;
        for(int c = 0; c < numChannels; c++) {
-               float cdiff = weightImage[c*channel_offset + y*w+x] - weightImage[c*channel_offset + (y+dy)*w+(x+dx)];
-               float pvar = varianceImage[c*channel_offset + y*w+x];
-               float qvar = varianceImage[c*channel_offset + (y+dy)*w+(x+dx)];
+               float cdiff = weight_image[c*channel_offset + y*w+x] - weight_image[c*channel_offset + (y+dy)*w+(x+dx)];
+               float pvar = variance_image[c*channel_offset + y*w+x];
+               float qvar = variance_image[c*channel_offset + (y+dy)*w+(x+dx)];
                diff += (cdiff*cdiff - a*(pvar + min(pvar, qvar))) / (1e-8f + k_2*(pvar+qvar));
        }
        if(numChannels > 1) {
                diff *= 1.0f/numChannels;
        }
-       differenceImage[y*w+x] = diff;
+       difference_image[y*w+x] = diff;
 }
 
 ccl_device_inline void kernel_filter_nlm_blur(int x, int y,
-                                              ccl_global float ccl_restrict_ptr differenceImage,
-                                              ccl_global float *outImage,
+                                              const ccl_global float *ccl_restrict difference_image,
+                                              ccl_global float *out_image,
                                               int4 rect, int w, int f)
 {
        float sum = 0.0f;
        const int low = max(rect.y, y-f);
        const int high = min(rect.w, y+f+1);
        for(int y1 = low; y1 < high; y1++) {
-               sum += differenceImage[y1*w+x];
+               sum += difference_image[y1*w+x];
        }
        sum *= 1.0f/(high-low);
-       outImage[y*w+x] = sum;
+       out_image[y*w+x] = sum;
 }
 
 ccl_device_inline void kernel_filter_nlm_calc_weight(int x, int y,
-                                                     ccl_global float ccl_restrict_ptr differenceImage,
-                                                     ccl_global float *outImage,
+                                                     const ccl_global float *ccl_restrict difference_image,
+                                                     ccl_global float *out_image,
                                                      int4 rect, int w, int f)
 {
        float sum = 0.0f;
        const int low = max(rect.x, x-f);
        const int high = min(rect.z, x+f+1);
        for(int x1 = low; x1 < high; x1++) {
-               sum += differenceImage[y*w+x1];
+               sum += difference_image[y*w+x1];
        }
        sum *= 1.0f/(high-low);
-       outImage[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,
                                                        int dx, int dy,
-                                                       ccl_global float ccl_restrict_ptr differenceImage,
-                                                       ccl_global float ccl_restrict_ptr image,
-                                                       ccl_global float *outImage,
-                                                       ccl_global float *accumImage,
+                                                       const ccl_global float *ccl_restrict difference_image,
+                                                       const ccl_global float *ccl_restrict image,
+                                                       ccl_global float *out_image,
+                                                       ccl_global float *accum_image,
                                                        int4 rect, int w, int f)
 {
        float sum = 0.0f;
        const int low = max(rect.x, x-f);
        const int high = min(rect.z, x+f+1);
        for(int x1 = low; x1 < high; x1++) {
-               sum += differenceImage[y*w+x1];
+               sum += difference_image[y*w+x1];
        }
        sum *= 1.0f/(high-low);
-       if(outImage) {
-               accumImage[y*w+x] += sum;
-               outImage[y*w+x] += sum*image[(y+dy)*w+(x+dx)];
+       if(out_image) {
+               accum_image[y*w+x] += sum;
+               out_image[y*w+x] += sum*image[(y+dy)*w+(x+dx)];
        }
        else {
-               accumImage[y*w+x] = sum;
+               accum_image[y*w+x] = sum;
        }
 }
 
 ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy,
                                                            int dx, int dy,
-                                                           ccl_global float ccl_restrict_ptr differenceImage,
-                                                           ccl_global float ccl_restrict_ptr buffer,
-                                                           ccl_global float *color_pass,
-                                                           ccl_global float *variance_pass,
-                                                           ccl_global float ccl_restrict_ptr transform,
+                                                           const ccl_global float *ccl_restrict difference_image,
+                                                           const ccl_global float *ccl_restrict buffer,
+                                                           const ccl_global float *ccl_restrict transform,
                                                            ccl_global int *rank,
                                                            ccl_global float *XtWX,
                                                            ccl_global float3 *XtWY,
@@ -115,7 +113,7 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy,
        const int high = min(rect.z, x+f+1);
        float sum = 0.0f;
        for(int x1 = low; x1 < high; x1++) {
-               sum += differenceImage[y*w+x1];
+               sum += difference_image[y*w+x1];
        }
        float weight = sum * (1.0f/(high - low));
 
@@ -130,18 +128,17 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy,
                                        dx, dy, w, h,
                                        pass_stride,
                                        buffer,
-                                       color_pass, variance_pass,
                                        transform, rank,
                                        weight, XtWX, XtWY,
                                        localIdx);
 }
 
 ccl_device_inline void kernel_filter_nlm_normalize(int x, int y,
-                                                   ccl_global float *outImage,
-                                                   ccl_global float ccl_restrict_ptr accumImage,
+                                                   ccl_global float *out_image,
+                                                   const ccl_global float *ccl_restrict accum_image,
                                                    int4 rect, int w)
 {
-       outImage[y*w+x] /= accumImage[y*w+x];
+       out_image[y*w+x] /= accum_image[y*w+x];
 }
 
 CCL_NAMESPACE_END
index 54bcf8880523626d130082e29740806d7755e50d..d5ae1b7392754269f643ee440e72b9b4887b5afd 100644 (file)
@@ -44,7 +44,7 @@ ccl_device void kernel_filter_divide_shadow(int sample,
 
        int offset = tiles->offsets[tile];
        int stride = tiles->strides[tile];
-       ccl_global float ccl_restrict_ptr center_buffer = (ccl_global float*) tiles->buffers[tile];
+       const ccl_global float *ccl_restrict center_buffer = (ccl_global float*) tiles->buffers[tile];
        center_buffer += (y*stride + x + offset)*buffer_pass_stride;
        center_buffer += buffer_denoising_offset + 14;
 
@@ -104,6 +104,66 @@ ccl_device void kernel_filter_get_feature(int sample,
        }
 }
 
+ccl_device void kernel_filter_detect_outliers(int x, int y,
+                                              ccl_global float *image,
+                                              ccl_global float *variance,
+                                              ccl_global float *depth,
+                                              ccl_global float *out,
+                                              int4 rect,
+                                              int pass_stride)
+{
+       int buffer_w = align_up(rect.z - rect.x, 4);
+
+       int n = 0;
+       float values[25];
+       for(int y1 = max(y-2, rect.y); y1 < min(y+3, rect.w); y1++) {
+               for(int x1 = max(x-2, rect.x); x1 < min(x+3, rect.z); x1++) {
+                       int idx = (y1-rect.y)*buffer_w + (x1-rect.x);
+                       float L = average(make_float3(image[idx], image[idx+pass_stride], image[idx+2*pass_stride]));
+
+                       /* Find the position of L. */
+                       int i;
+                       for(i = 0; i < n; i++) {
+                               if(values[i] > L) break;
+                       }
+                       /* Make space for L by shifting all following values to the right. */
+                       for(int j = n; j > i; j--) {
+                               values[j] = values[j-1];
+                       }
+                       /* Insert L. */
+                       values[i] = L;
+                       n++;
+               }
+       }
+
+       int idx = (y-rect.y)*buffer_w + (x-rect.x);
+       float L = average(make_float3(image[idx], image[idx+pass_stride], image[idx+2*pass_stride]));
+
+       float ref = 2.0f*values[(int)(n*0.75f)];
+       float fac = 1.0f;
+       if(L > ref) {
+               /* The pixel appears to be an outlier.
+                * However, it may just be a legitimate highlight. Therefore, it is checked how likely it is that the pixel
+                * should actually be at the reference value:
+                * If the reference is within the 3-sigma interval, the pixel is assumed to be a statistical outlier.
+                * Otherwise, it is very unlikely that the pixel should be darker, which indicates a legitimate highlight.
+                */
+               float stddev = sqrtf(average(make_float3(variance[idx], variance[idx+pass_stride], variance[idx+2*pass_stride])));
+               if(L - 3*stddev < ref) {
+                       /* The pixel is an outlier, so negate the depth value to mark it as one.
+                        * Also, scale its brightness down to the outlier threshold to avoid trouble with the NLM weights. */
+                       depth[idx] = -depth[idx];
+                       fac = ref/L;
+                       variance[idx              ] *= fac*fac;
+                       variance[idx + pass_stride] *= fac*fac;
+                       variance[idx+2*pass_stride] *= fac*fac;
+               }
+       }
+       out[idx              ] = fac*image[idx];
+       out[idx + pass_stride] = fac*image[idx + pass_stride];
+       out[idx+2*pass_stride] = fac*image[idx+2*pass_stride];
+}
+
 /* Combine A/B buffers.
  * Calculates the combined mean and the buffer variance. */
 ccl_device void kernel_filter_combine_halves(int x, int y,
index 02f3802fa0cda2897c11323662d2a77d806f6f17..25a3025056c0bc0469cb5d90ee2498ca0beef390 100644 (file)
@@ -21,40 +21,39 @@ ccl_device_inline void kernel_filter_construct_gramian(int x, int y,
                                                        int dx, int dy,
                                                        int w, int h,
                                                        int pass_stride,
-                                                       ccl_global float ccl_restrict_ptr buffer,
-                                                       ccl_global float *color_pass,
-                                                       ccl_global float *variance_pass,
-                                                       ccl_global float ccl_restrict_ptr transform,
+                                                       const ccl_global float *ccl_restrict buffer,
+                                                       const ccl_global float *ccl_restrict transform,
                                                        ccl_global int *rank,
                                                        float weight,
                                                        ccl_global float *XtWX,
                                                        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
 
-       float3 p_color = filter_get_pixel_color(color_pass + p_offset, pass_stride);
-       float3 q_color = filter_get_pixel_color(color_pass + q_offset, pass_stride);
-
-       float p_std_dev = sqrtf(filter_get_pixel_variance(variance_pass + p_offset, pass_stride));
-       float q_std_dev = sqrtf(filter_get_pixel_variance(variance_pass + q_offset, pass_stride));
+       float3 q_color = filter_get_color(buffer + q_offset, pass_stride);
 
-       if(average(fabs(p_color - q_color)) > 3.0f*(p_std_dev + q_std_dev + 1e-3f)) {
+       /* If the pixel was flagged as an outlier during prefiltering, skip it. */
+       if(ccl_get_feature(buffer + q_offset, 0) < 0.0f) {
                return;
        }
 
@@ -75,16 +74,33 @@ 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. */
+       float3 mean_color = XtWY[0]/XtWX[0];
+
        math_trimatrix_vec3_solve(XtWX, XtWY, (*rank)+1, stride);
 
        float3 final_color = XtWY[0];
+       if(!isfinite3_safe(final_color)) {
+               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;
@@ -98,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
index 139dc402d216f81d72096a292a8b4ee783089589..a5f87c05ec0b89ea20e1c886a916d132db8d6961 100644 (file)
@@ -16,7 +16,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
+ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buffer,
                                                   int x, int y, int4 rect,
                                                   int pass_stride,
                                                   float *transform, int *rank,
@@ -29,20 +29,15 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
        /* Temporary storage, used in different steps of the algorithm. */
        float tempmatrix[DENOISE_FEATURES*DENOISE_FEATURES];
        float tempvector[2*DENOISE_FEATURES];
-       float ccl_restrict_ptr pixel_buffer;
+       const float *ccl_restrict pixel_buffer;
        int2 pixel;
 
-
-
-
        /* === Calculate denoising window. === */
        int2 low  = make_int2(max(rect.x, x - radius),
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
-
-
-
+       int num_pixels = (high.y - low.y) * (high.x - low.x);
 
        /* === Shift feature passes to have mean 0. === */
        float feature_means[DENOISE_FEATURES];
@@ -52,8 +47,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                math_vector_add(feature_means, features, DENOISE_FEATURES);
        } END_FOR_PIXEL_WINDOW
 
-       float pixel_scale = 1.0f / ((high.y - low.y) * (high.x - low.x));
-       math_vector_scale(feature_means, pixel_scale, DENOISE_FEATURES);
+       math_vector_scale(feature_means, 1.0f / num_pixels, DENOISE_FEATURES);
 
        /* === Scale the shifted feature passes to a range of [-1; 1], will be baked into the transform later. === */
        float *feature_scale = tempvector;
@@ -66,7 +60,6 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
 
        filter_calculate_scale(feature_scale);
 
-
        /* === Generate the feature transformation. ===
         * This transformation maps the DENOISE_FEATURES-dimentional feature space to a reduced feature (r-feature) space
         * which generally has fewer dimensions. This mainly helps to prevent overfitting. */
@@ -80,6 +73,8 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
 
        math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, 1);
        *rank = 0;
+       /* Prevent overfitting when a small window is used. */
+       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
                for(int i = 0; i < DENOISE_FEATURES; i++) {
@@ -88,7 +83,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                threshold_energy *= 1.0f - (-pca_threshold);
 
                float reduced_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
@@ -96,7 +91,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                }
        }
        else {
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
index 68304e14143a2ecbb05289554f0723c4171e09f2..83a1222bbdb25d2818bd1c4d63c4654f2830b145 100644 (file)
@@ -16,7 +16,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_ptr buffer,
+ccl_device void kernel_filter_construct_transform(const ccl_global float *ccl_restrict buffer,
                                                   int x, int y, int4 rect,
                                                   int pass_stride,
                                                   ccl_global float *transform,
@@ -38,7 +38,8 @@ ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
-       ccl_global float ccl_restrict_ptr pixel_buffer;
+       int num_pixels = (high.y - low.y) * (high.x - low.x);
+       const ccl_global float *ccl_restrict pixel_buffer;
        int2 pixel;
 
 
@@ -52,8 +53,7 @@ ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_
                math_vector_add(feature_means, features, DENOISE_FEATURES);
        } END_FOR_PIXEL_WINDOW
 
-       float pixel_scale = 1.0f / ((high.y - low.y) * (high.x - low.x));
-       math_vector_scale(feature_means, pixel_scale, DENOISE_FEATURES);
+       math_vector_scale(feature_means, 1.0f / num_pixels, DENOISE_FEATURES);
 
        /* === Scale the shifted feature passes to a range of [-1; 1], will be baked into the transform later. === */
        float feature_scale[DENOISE_FEATURES];
@@ -81,6 +81,8 @@ ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_
 
        math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, transform_stride);
        *rank = 0;
+       /* Prevent overfitting when a small window is used. */
+       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
                for(int i = 0; i < DENOISE_FEATURES; i++) {
@@ -89,7 +91,7 @@ ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_
                threshold_energy *= 1.0f - (-pca_threshold);
 
                float reduced_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
@@ -97,7 +99,7 @@ ccl_device void kernel_filter_construct_transform(ccl_global float ccl_restrict_
                }
        }
        else {
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
index ed3a92f6241611c07699acf16770d193bbcf8fc6..30dc2969b114f40fd3c41a360e9a5ad0ab4efaa9 100644 (file)
@@ -16,7 +16,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
+ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buffer,
                                                   int x, int y, int4 rect,
                                                   int pass_stride,
                                                   float *transform, int *rank,
@@ -25,13 +25,14 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
        int buffer_w = align_up(rect.z - rect.x, 4);
 
        __m128 features[DENOISE_FEATURES];
-       float ccl_restrict_ptr pixel_buffer;
+       const float *ccl_restrict pixel_buffer;
        int2 pixel;
 
        int2 low  = make_int2(max(rect.x, x - radius),
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
+       int num_pixels = (high.y - low.y) * (high.x - low.x);
 
        __m128 feature_means[DENOISE_FEATURES];
        math_vector_zero_sse(feature_means, DENOISE_FEATURES);
@@ -40,7 +41,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                math_vector_add_sse(feature_means, DENOISE_FEATURES, features);
        } END_FOR_PIXEL_WINDOW_SSE
 
-       __m128 pixel_scale = _mm_set1_ps(1.0f / ((high.y - low.y) * (high.x - low.x)));
+       __m128 pixel_scale = _mm_set1_ps(1.0f / num_pixels);
        for(int i = 0; i < DENOISE_FEATURES; i++) {
                feature_means[i] = _mm_mul_ps(_mm_hsum_ps(feature_means[i]), pixel_scale);
        }
@@ -68,6 +69,8 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
        math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, 1);
 
        *rank = 0;
+       /* Prevent overfitting when a small window is used. */
+       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
                for(int i = 0; i < DENOISE_FEATURES; i++) {
@@ -76,7 +79,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                threshold_energy *= 1.0f - (-pca_threshold);
 
                float reduced_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
@@ -84,7 +87,7 @@ ccl_device void kernel_filter_construct_transform(float ccl_restrict_ptr buffer,
                }
        }
        else {
-               for(int i = 0; i < DENOISE_FEATURES; i++, (*rank)++) {
+               for(int i = 0; i < max_rank; i++, (*rank)++) {
                        float s = feature_matrix[i*DENOISE_FEATURES+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
index 06728415c155a8039131bd08a509d83529d9dae5..175bd6b9737daf999c6519b622781fe615a97639 100644 (file)
@@ -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__
index 7595e74e2d569873541d23eefe2a87b714a14ace..93934ee6b380aefee4930425333d239d9e282bd9 100644 (file)
@@ -42,8 +42,6 @@
 #include "util/util_types.h"
 #include "util/util_texture.h"
 
-#define ccl_restrict_ptr const * __restrict
-
 #define ccl_addr_space
 
 #define ccl_local_id(d) 0
@@ -197,7 +195,7 @@ template<typename T> struct texture_image  {
                                        if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
                                                return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                        }
-                                       /* Fall through. */
+                                       ATTR_FALLTHROUGH;
                                case EXTENSION_EXTEND:
                                        ix = wrap_clamp(ix, width);
                                        iy = wrap_clamp(iy, height);
@@ -224,7 +222,7 @@ template<typename T> struct texture_image  {
                                        if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
                                                return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                        }
-                                       /* Fall through. */
+                                       ATTR_FALLTHROUGH;
                                case EXTENSION_EXTEND:
                                        nix = wrap_clamp(ix+1, width);
                                        niy = wrap_clamp(iy+1, height);
@@ -267,7 +265,7 @@ template<typename T> struct texture_image  {
                                        if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
                                                return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                        }
-                                       /* Fall through. */
+                                       ATTR_FALLTHROUGH;
                                case EXTENSION_EXTEND:
                                        pix = wrap_clamp(ix-1, width);
                                        piy = wrap_clamp(iy-1, height);
@@ -337,7 +335,7 @@ template<typename T> struct texture_image  {
                                {
                                        return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                }
-                               /* Fall through. */
+                               ATTR_FALLTHROUGH;
                        case EXTENSION_EXTEND:
                                ix = wrap_clamp(ix, width);
                                iy = wrap_clamp(iy, height);
@@ -376,7 +374,7 @@ template<typename T> struct texture_image  {
                                {
                                        return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                }
-                               /* Fall through. */
+                               ATTR_FALLTHROUGH;
                        case EXTENSION_EXTEND:
                                nix = wrap_clamp(ix+1, width);
                                niy = wrap_clamp(iy+1, height);
@@ -451,7 +449,7 @@ template<typename T> struct texture_image  {
                                {
                                        return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                                }
-                               /* Fall through. */
+                               ATTR_FALLTHROUGH;
                        case EXTENSION_EXTEND:
                                pix = wrap_clamp(ix-1, width);
                                piy = wrap_clamp(iy-1, height);
index 80d7401fbcfefe6de44709bb025a1aacc1fa0bbb..38708f7ff0b77910ed634e0470637305ea8a393e 100644 (file)
@@ -55,7 +55,8 @@
 #define ccl_restrict __restrict__
 #define ccl_align(n) __align__(n)
 
-#define ccl_restrict_ptr const * __restrict__
+#define ATTR_FALLTHROUGH
+
 #define CCL_MAX_LOCAL_SIZE (CUDA_THREADS_BLOCK_WIDTH*CUDA_THREADS_BLOCK_WIDTH)
 
 
index 15cf4b81b21a210d5886c46bf68f3708387d1d65..4836c2903120b8b78f4989de9390546bcee4b13e 100644 (file)
@@ -50,7 +50,7 @@
 #  define ccl_addr_space
 #endif
 
-#define ccl_restrict_ptr const * __restrict__
+#define ATTR_FALLTHROUGH
 
 #define ccl_local_id(d) get_local_id(d)
 #define ccl_global_id(d) get_global_id(d)
index c694e60ce143f62ff5ca47c65db64e48b59052d7..9d52834ffcc7c7c627e225d67c4f7ae564643a5a 100644 (file)
@@ -152,6 +152,11 @@ ccl_device_inline void kernel_update_denoising_features(KernelGlobals *kg,
 
        L->denoising_depth += ensure_finite(state->denoising_feature_weight * sd->ray_length);
 
+       /* Skip implicitly transparent surfaces. */
+       if(sd->flag & SD_HAS_ONLY_VOLUME) {
+               return;
+       }
+
        float3 normal = make_float3(0.0f, 0.0f, 0.0f);
        float3 albedo = make_float3(0.0f, 0.0f, 0.0f);
        float sum_weight = 0.0f, sum_nonspecular_weight = 0.0f;
@@ -356,7 +361,16 @@ ccl_device_inline void kernel_write_result(KernelGlobals *kg, ccl_global float *
 #  endif
                        if(kernel_data.film.pass_denoising_clean) {
                                float3 noisy, clean;
-                               path_radiance_split_denoising(kg, L, &noisy, &clean);
+#ifdef __SHADOW_TRICKS__
+                               if(is_shadow_catcher) {
+                                       noisy = L_sum;
+                                       clean = make_float3(0.0f, 0.0f, 0.0f);
+                               }
+                               else
+#endif  /* __SHADOW_TRICKS__ */
+                               {
+                                       path_radiance_split_denoising(kg, L, &noisy, &clean);
+                               }
                                kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR,
                                                                  sample, noisy);
                                kernel_write_pass_float3_unaligned(buffer + kernel_data.film.pass_denoising_clean,
index 0fa77d9e8bd3d215ee612ef9ee4e025c16d56d5b..5d92fd1220177bedf54a461e77c60447480e0f71 100644 (file)
@@ -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)
index 96bc636d5acaca0608cf55dc07de35ac6761e1e1..e32d4bbbc1b423154aa5a19fc50fa020709b1173 100644 (file)
@@ -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__
index dbeaffdfb24975b9ad2ee239bd96f49d16ce38a9..31e47e837fda4b5b184deb674cfaced1fdd9c79a 100644 (file)
@@ -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))
index 10007ee2635870fe7a04031bb8dd3647a06dc187..2ed713299fd88d526673fd60deb2e7b0296afb92 100644 (file)
@@ -43,6 +43,14 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
                                                    int buffer_denoising_offset,
                                                    bool use_split_variance);
 
+void KERNEL_FUNCTION_FULL_NAME(filter_detect_outliers)(int x, int y,
+                                                       ccl_global float *image,
+                                                       ccl_global float *variance,
+                                                       ccl_global float *depth,
+                                                       ccl_global float *output,
+                                                       int *rect,
+                                                       int pass_stride);
+
 void KERNEL_FUNCTION_FULL_NAME(filter_combine_halves)(int x, int y,
                                                       float *mean,
                                                       float *variance,
@@ -64,43 +72,41 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx,
                                                            int dy,
-                                                           float *weightImage,
+                                                           float *weight_image,
                                                            float *variance,
-                                                           float *differenceImage,
+                                                           float *difference_image,
                                                            int* rect,
                                                            int w,
                                                            int channel_offset,
                                                            float a,
                                                            float k_2);
 
-void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *differenceImage,
-                                                float *outImage,
+void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *difference_image,
+                                                float *out_image,
                                                 int* rect,
                                                 int w,
                                                 int f);
 
-void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *differenceImage,
-                                                       float *outImage,
+void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *difference_image,
+                                                       float *out_image,
                                                        int* rect,
                                                        int w,
                                                        int f);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx,
                                                          int dy,
-                                                         float *differenceImage,
+                                                         float *difference_image,
                                                          float *image,
-                                                         float *outImage,
-                                                         float *accumImage,
+                                                         float *out_image,
+                                                         float *accum_image,
                                                          int* rect,
                                                          int w,
                                                          int f);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int dy,
-                                                             float *differenceImage,
+                                                             float *difference_image,
                                                              float *buffer,
-                                                             float *color_pass,
-                                                             float *variance_pass,
                                                              float *transform,
                                                              int *rank,
                                                              float *XtWX,
@@ -112,8 +118,8 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int f,
                                                              int pass_stride);
 
-void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *outImage,
-                                                     float *accumImage,
+void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *out_image,
+                                                     float *accum_image,
                                                      int* rect,