archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity

Merge branch 'blender2.8' into gpencil_2.8

Browse Source
This commit is contained in:
Antonio Vazquez
2017-03-10 11:10:13 +01:00
parent 8bfd2a83ff 6a8cb92404
commit 51395e01c3
68 changed files with 1950 additions and 929 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
intern/cycles/blender/addon/properties.py
View File

@@ -695,6 +695,8 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
update=devices_update_callback
)
cls.debug_opencl_kernel_single_program = BoolProperty(name="Single Program", default=False, update=devices_update_callback);
cls.debug_use_opencl_debug = BoolProperty(name="Debug OpenCL", default=False)
@classmethod

1
intern/cycles/blender/addon/ui.py
View File

@@ -1529,6 +1529,7 @@ class CyclesRender_PT_debug(CyclesButtonsPanel, Panel):
col.label('OpenCL Flags:')
col.prop(cscene, "debug_opencl_kernel_type", text="Kernel")
col.prop(cscene, "debug_opencl_device_type", text="Device")
col.prop(cscene, "debug_opencl_kernel_single_program", text="Single Program")
col.prop(cscene, "debug_use_opencl_debug", text="Debug")

1
intern/cycles/blender/blender_python.cpp
View File

@@ -106,6 +106,7 @@ bool debug_flags_sync_from_scene(BL::Scene b_scene)
}
/* Synchronize other OpenCL flags. */
flags.opencl.debug = get_boolean(cscene, "debug_use_opencl_debug");
flags.opencl.single_program = get_boolean(cscene, "debug_opencl_kernel_single_program");
return flags.opencl.device_type != opencl_device_type ||
flags.opencl.kernel_type != opencl_kernel_type;
}

2
intern/cycles/device/device.h
View File

@@ -194,7 +194,7 @@ public:
if(!use_patch_evaluation) {
build_options += " -D__NO_PATCH_EVAL__";
}
if(!use_transparent) {
if(!use_transparent && !use_volume) {
build_options += " -D__NO_TRANSPARENT__";
}
return build_options;

43
intern/cycles/device/device_split_kernel.cpp
View File

@@ -35,13 +35,18 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device) : device(device)
kernel_path_init = NULL;
kernel_scene_intersect = NULL;
kernel_lamp_emission = NULL;
kernel_do_volume = NULL;
kernel_queue_enqueue = NULL;
kernel_background_buffer_update = NULL;
kernel_indirect_background = NULL;
kernel_shader_eval = NULL;
kernel_holdout_emission_blurring_pathtermination_ao = NULL;
kernel_subsurface_scatter = NULL;
kernel_direct_lighting = NULL;
kernel_shadow_blocked = NULL;
kernel_shadow_blocked_ao = NULL;
kernel_shadow_blocked_dl = NULL;
kernel_next_iteration_setup = NULL;
kernel_indirect_subsurface = NULL;
kernel_buffer_update = NULL;
}
DeviceSplitKernel::~DeviceSplitKernel()
@@ -55,13 +60,18 @@ DeviceSplitKernel::~DeviceSplitKernel()
delete kernel_path_init;
delete kernel_scene_intersect;
delete kernel_lamp_emission;
delete kernel_do_volume;
delete kernel_queue_enqueue;
delete kernel_background_buffer_update;
delete kernel_indirect_background;
delete kernel_shader_eval;
delete kernel_holdout_emission_blurring_pathtermination_ao;
delete kernel_subsurface_scatter;
delete kernel_direct_lighting;
delete kernel_shadow_blocked;
delete kernel_shadow_blocked_ao;
delete kernel_shadow_blocked_dl;
delete kernel_next_iteration_setup;
delete kernel_indirect_subsurface;
delete kernel_buffer_update;
}
bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures& requested_features)
@@ -75,13 +85,18 @@ bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures& requested_fe
LOAD_KERNEL(path_init);
LOAD_KERNEL(scene_intersect);
LOAD_KERNEL(lamp_emission);
LOAD_KERNEL(do_volume);
LOAD_KERNEL(queue_enqueue);
LOAD_KERNEL(background_buffer_update);
LOAD_KERNEL(indirect_background);
LOAD_KERNEL(shader_eval);
LOAD_KERNEL(holdout_emission_blurring_pathtermination_ao);
LOAD_KERNEL(subsurface_scatter);
LOAD_KERNEL(direct_lighting);
LOAD_KERNEL(shadow_blocked);
LOAD_KERNEL(shadow_blocked_ao);
LOAD_KERNEL(shadow_blocked_dl);
LOAD_KERNEL(next_iteration_setup);
LOAD_KERNEL(indirect_subsurface);
LOAD_KERNEL(buffer_update);
#undef LOAD_KERNEL
@@ -210,23 +225,23 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
bool activeRaysAvailable = true;
while(activeRaysAvailable) {
/* Twice the global work size of other kernels for
* ckPathTraceKernel_shadow_blocked_direct_lighting. */
size_t global_size_shadow_blocked[2];
global_size_shadow_blocked[0] = global_size[0] * 2;
global_size_shadow_blocked[1] = global_size[1];
/* Do path-iteration in host [Enqueue Path-iteration kernels. */
for(int PathIter = 0; PathIter < 16; PathIter++) {
ENQUEUE_SPLIT_KERNEL(scene_intersect, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(lamp_emission, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(do_volume, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(background_buffer_update, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(indirect_background, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shader_eval, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(holdout_emission_blurring_pathtermination_ao, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(subsurface_scatter, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shadow_blocked, global_size_shadow_blocked, 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(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);
ENQUEUE_SPLIT_KERNEL(buffer_update, global_size, local_size);
if(task->get_cancel()) {
return true;

9
intern/cycles/device/device_split_kernel.h
View File

@@ -58,13 +58,18 @@ private:
SplitKernelFunction *kernel_path_init;
SplitKernelFunction *kernel_scene_intersect;
SplitKernelFunction *kernel_lamp_emission;
SplitKernelFunction *kernel_do_volume;
SplitKernelFunction *kernel_queue_enqueue;
SplitKernelFunction *kernel_background_buffer_update;
SplitKernelFunction *kernel_indirect_background;
SplitKernelFunction *kernel_shader_eval;
SplitKernelFunction *kernel_holdout_emission_blurring_pathtermination_ao;
SplitKernelFunction *kernel_subsurface_scatter;
SplitKernelFunction *kernel_direct_lighting;
SplitKernelFunction *kernel_shadow_blocked;
SplitKernelFunction *kernel_shadow_blocked_ao;
SplitKernelFunction *kernel_shadow_blocked_dl;
SplitKernelFunction *kernel_next_iteration_setup;
SplitKernelFunction *kernel_indirect_subsurface;
SplitKernelFunction *kernel_buffer_update;
/* Global memory variables [porting]; These memory is used for
* co-operation between different kernels; Data written by one

1
intern/cycles/device/opencl/opencl.h
View File

@@ -90,6 +90,7 @@ public:
cl_device_id device_id);
static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
bool force_all = false);
static bool use_single_program();
};
/* Thread safe cache for contexts and programs.

21
intern/cycles/device/opencl/opencl_split.cpp
View File

@@ -77,16 +77,18 @@ public:
virtual bool load_kernels(const DeviceRequestedFeatures& requested_features,
vector<OpenCLDeviceBase::OpenCLProgram*> &programs)
{
bool single_program = OpenCLInfo::use_single_program();
program_data_init = OpenCLDeviceBase::OpenCLProgram(this,
"split_data_init",
"kernel_data_init.cl",
single_program ? "split" : "split_data_init",
single_program ? "kernel_split.cl" : "kernel_data_init.cl",
get_build_options(this, requested_features));
program_data_init.add_kernel(ustring("path_trace_data_init"));
programs.push_back(&program_data_init);
program_state_buffer_size = OpenCLDeviceBase::OpenCLProgram(this,
"split_state_buffer_size",
"kernel_state_buffer_size.cl",
single_program ? "split" : "split_state_buffer_size",
single_program ? "kernel_split.cl" : "kernel_state_buffer_size.cl",
get_build_options(this, requested_features));
program_state_buffer_size.add_kernel(ustring("path_trace_state_buffer_size"));
programs.push_back(&program_state_buffer_size);
@@ -207,10 +209,13 @@ public:
{
OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device);
kernel->program = OpenCLDeviceBase::OpenCLProgram(device,
"split_" + kernel_name,
"kernel_" + kernel_name + ".cl",
get_build_options(device, requested_features));
bool single_program = OpenCLInfo::use_single_program();
kernel->program =
OpenCLDeviceBase::OpenCLProgram(device,
single_program ? "split" : "split_" + kernel_name,
single_program ? "kernel_split.cl" : "kernel_" + kernel_name + ".cl",
get_build_options(device, requested_features));
kernel->program.add_kernel(ustring("path_trace_" + kernel_name));
kernel->program.load();

5
intern/cycles/device/opencl/opencl_util.cpp
View File

@@ -552,6 +552,11 @@ bool OpenCLInfo::use_debug()
return DebugFlags().opencl.debug;
}
bool OpenCLInfo::use_single_program()
{
return DebugFlags().opencl.single_program;
}
bool OpenCLInfo::kernel_use_advanced_shading(const string& platform)
{
/* keep this in sync with kernel_types.h! */

30
intern/cycles/kernel/CMakeLists.txt
View File

@@ -16,17 +16,23 @@ set(SRC
kernels/cpu/kernel_split.cpp
kernels/opencl/kernel.cl
kernels/opencl/kernel_state_buffer_size.cl
kernels/opencl/kernel_split.cl
kernels/opencl/kernel_data_init.cl
kernels/opencl/kernel_path_init.cl
kernels/opencl/kernel_queue_enqueue.cl
kernels/opencl/kernel_scene_intersect.cl
kernels/opencl/kernel_lamp_emission.cl
kernels/opencl/kernel_background_buffer_update.cl
kernels/opencl/kernel_do_volume.cl
kernels/opencl/kernel_indirect_background.cl
kernels/opencl/kernel_shader_eval.cl
kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
kernels/opencl/kernel_subsurface_scatter.cl
kernels/opencl/kernel_direct_lighting.cl
kernels/opencl/kernel_shadow_blocked.cl
kernels/opencl/kernel_shadow_blocked_ao.cl
kernels/opencl/kernel_shadow_blocked_dl.cl
kernels/opencl/kernel_next_iteration_setup.cl
kernels/opencl/kernel_indirect_subsurface.cl
kernels/opencl/kernel_buffer_update.cl
kernels/cuda/kernel.cu
kernels/cuda/kernel_split.cu
)
@@ -71,6 +77,7 @@ set(SRC_HEADERS
kernel_path_common.h
kernel_path_state.h
kernel_path_surface.h
kernel_path_subsurface.h
kernel_path_volume.h
kernel_projection.h
kernel_queues.h
@@ -196,20 +203,25 @@ set(SRC_UTIL_HEADERS
)
set(SRC_SPLIT_HEADERS
split/kernel_background_buffer_update.h
split/kernel_buffer_update.h
split/kernel_data_init.h
split/kernel_direct_lighting.h
split/kernel_do_volume.h
split/kernel_holdout_emission_blurring_pathtermination_ao.h
split/kernel_indirect_background.h
split/kernel_indirect_subsurface.h
split/kernel_lamp_emission.h
split/kernel_next_iteration_setup.h
split/kernel_path_init.h
split/kernel_queue_enqueue.h
split/kernel_scene_intersect.h
split/kernel_shader_eval.h
split/kernel_shadow_blocked.h
split/kernel_shadow_blocked_ao.h
split/kernel_shadow_blocked_dl.h
split/kernel_split_common.h
split/kernel_split_data.h
split/kernel_split_data_types.h
split/kernel_subsurface_scatter.h
)
# CUDA module
@@ -402,17 +414,23 @@ endif()
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_state_buffer_size.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_split.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_data_init.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_path_init.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_queue_enqueue.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_scene_intersect.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_lamp_emission.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_background_buffer_update.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_do_volume.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_indirect_background.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shader_eval.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_subsurface_scatter.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_direct_lighting.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_ao.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_dl.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_next_iteration_setup.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_indirect_subsurface.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_buffer_update.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel.cu" ${CYCLES_INSTALL_PATH}/kernel/kernels/cuda)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel_split.cu" ${CYCLES_INSTALL_PATH}/kernel/kernels/cuda)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_HEADERS}" ${CYCLES_INSTALL_PATH}/kernel)

10
intern/cycles/kernel/bvh/bvh_shadow_all.h
View File

@@ -309,9 +309,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
object = kernel_tex_fetch(__prim_object, -prim_addr-1);
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect_t, &ob_itfm);
isect_t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect_t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
@@ -362,12 +362,10 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
}
}
else {
float ignore_t = FLT_MAX;
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &ignore_t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &ignore_t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
}

18
intern/cycles/kernel/bvh/bvh_subsurface.h
View File

@@ -75,16 +75,16 @@ void BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
if(!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
bvh_instance_motion_push(kg,
subsurface_object,
ray,
&P,
&dir,
&idir,
&isect_t,
&ob_itfm);
isect_t = bvh_instance_motion_push(kg,
subsurface_object,
ray,
&P,
&dir,
&idir,
isect_t,
&ob_itfm);
#else
bvh_instance_push(kg, subsurface_object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, subsurface_object, ray, &P, &dir, &idir, isect_t);
#endif
object = subsurface_object;
}

8
intern/cycles/kernel/bvh/bvh_traversal.h
View File

@@ -354,9 +354,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
object = kernel_tex_fetch(__prim_object, -prim_addr-1);
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
@@ -391,9 +391,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
/* instance pop */
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);

8
intern/cycles/kernel/bvh/bvh_volume.h
View File

@@ -238,9 +238,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
int object_flag = kernel_tex_fetch(__object_flag, object);
if(object_flag & SD_OBJECT_HAS_VOLUME) {
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
@@ -281,9 +281,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
/* instance pop */
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);

10
intern/cycles/kernel/bvh/bvh_volume_all.h
View File

@@ -288,11 +288,10 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
object = kernel_tex_fetch(__prim_object, -prim_addr-1);
int object_flag = kernel_tex_fetch(__object_flag, object);
if(object_flag & SD_OBJECT_HAS_VOLUME) {
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect_t, &ob_itfm);
isect_t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect_t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
@@ -348,11 +347,10 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
}
}
else {
float ignore_t = FLT_MAX;
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &ignore_t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &ignore_t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
triangle_intersect_precalc(dir, &isect_precalc);
}

9
intern/cycles/kernel/bvh/qbvh_shadow_all.h
View File

@@ -390,9 +390,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
object = kernel_tex_fetch(__prim_object, -prim_addr-1);
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect_t, &ob_itfm);
isect_t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect_t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
num_hits_in_instance = 0;
@@ -445,11 +445,10 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
}
}
else {
float ignore_t = FLT_MAX;
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &ignore_t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &ignore_t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
}

18
intern/cycles/kernel/bvh/qbvh_subsurface.h
View File

@@ -64,16 +64,16 @@ ccl_device void BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
if(!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
bvh_instance_motion_push(kg,
subsurface_object,
ray,
&P,
&dir,
&idir,
&isect_t,
&ob_itfm);
isect_t = bvh_instance_motion_push(kg,
subsurface_object,
ray,
&P,
&dir,
&idir,
isect_t,
&ob_itfm);
#else
bvh_instance_push(kg, subsurface_object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, subsurface_object, ray, &P, &dir, &idir, isect_t);
#endif
object = subsurface_object;
}

4
intern/cycles/kernel/bvh/qbvh_traversal.h
View File

@@ -468,9 +468,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
/* Instance pop. */
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
qbvh_near_far_idx_calc(idir,

8
intern/cycles/kernel/bvh/qbvh_volume.h
View File

@@ -295,9 +295,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
int object_flag = kernel_tex_fetch(__object_flag, object);
if(object_flag & SD_OBJECT_HAS_VOLUME) {
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
qbvh_near_far_idx_calc(idir,
@@ -341,9 +341,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
/* Instance pop. */
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &isect->t);
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
qbvh_near_far_idx_calc(idir,

9
intern/cycles/kernel/bvh/qbvh_volume_all.h
View File

@@ -346,9 +346,9 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
int object_flag = kernel_tex_fetch(__object_flag, object);
if(object_flag & SD_OBJECT_HAS_VOLUME) {
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect_t, &ob_itfm);
isect_t = bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, isect_t, &ob_itfm);
# else
bvh_instance_push(kg, object, ray, &P, &dir, &idir, &isect_t);
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
qbvh_near_far_idx_calc(idir,
@@ -406,11 +406,10 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
}
}
else {
float ignore_t = FLT_MAX;
# if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, &ignore_t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, &ignore_t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
}

62
intern/cycles/kernel/geom/geom_object.h
View File

@@ -425,7 +425,13 @@ ccl_device_inline float3 bvh_inverse_direction(float3 dir)
/* Transform ray into object space to enter static object in BVH */
ccl_device_inline void bvh_instance_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, ccl_addr_space float *t)
ccl_device_inline float bvh_instance_push(KernelGlobals *kg,
int object,
const Ray *ray,
float3 *P,
float3 *dir,
float3 *idir,
float t)
{
Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
@@ -435,8 +441,11 @@ ccl_device_inline void bvh_instance_push(KernelGlobals *kg, int object, const Ra
*dir = bvh_clamp_direction(normalize_len(transform_direction(&tfm, ray->D), &len));
*idir = bvh_inverse_direction(*dir);
if(*t != FLT_MAX)
*t *= len;
if(t != FLT_MAX) {
t *= len;
}
return t;
}
#ifdef __QBVH__
@@ -473,16 +482,24 @@ ccl_device_inline void qbvh_instance_push(KernelGlobals *kg,
/* Transorm ray to exit static object in BVH */
ccl_device_inline void bvh_instance_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, ccl_addr_space float *t)
ccl_device_inline float bvh_instance_pop(KernelGlobals *kg,
int object,
const Ray *ray,
float3 *P,
float3 *dir,
float3 *idir,
float t)
{
if(*t != FLT_MAX) {
if(t != FLT_MAX) {
Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
*t /= len(transform_direction(&tfm, ray->D));
t /= len(transform_direction(&tfm, ray->D));
}
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
return t;
}
/* Same as above, but returns scale factor to apply to multiple intersection distances */
@@ -501,13 +518,13 @@ ccl_device_inline void bvh_instance_pop_factor(KernelGlobals *kg, int object, co
#ifdef __OBJECT_MOTION__
/* Transform ray into object space to enter motion blurred object in BVH */
ccl_device_inline void bvh_instance_motion_push(KernelGlobals *kg,
ccl_device_inline float bvh_instance_motion_push(KernelGlobals *kg,
int object,
const Ray *ray,
float3 *P,
float3 *dir,
float3 *idir,
ccl_addr_space float *t,
float t,
Transform *itfm)
{
object_fetch_transform_motion_test(kg, object, ray->time, itfm);
@@ -518,8 +535,11 @@ ccl_device_inline void bvh_instance_motion_push(KernelGlobals *kg,
*dir = bvh_clamp_direction(normalize_len(transform_direction(itfm, ray->D), &len));
*idir = bvh_inverse_direction(*dir);
if(*t != FLT_MAX)
*t *= len;
if(t != FLT_MAX) {
t *= len;
}
return t;
}
#ifdef __QBVH__
@@ -557,22 +577,24 @@ ccl_device_inline void qbvh_instance_motion_push(KernelGlobals *kg,
/* Transorm ray to exit motion blurred object in BVH */
ccl_device_inline void bvh_instance_motion_pop(KernelGlobals *kg,
int object,
const Ray *ray,
float3 *P,
float3 *dir,
float3 *idir,
ccl_addr_space float *t,
Transform *itfm)
ccl_device_inline float bvh_instance_motion_pop(KernelGlobals *kg,
int object,
const Ray *ray,
float3 *P,
float3 *dir,
float3 *idir,
float t,
Transform *itfm)
{
if(*t != FLT_MAX) {
*t /= len(transform_direction(itfm, ray->D));
if(t != FLT_MAX) {
t /= len(transform_direction(itfm, ray->D));
}
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
return t;
}
/* Same as above, but returns scale factor to apply to multiple intersection distances */

167
intern/cycles/kernel/kernel_path.h
View File

@@ -46,6 +46,7 @@
#include "kernel_path_common.h"
#include "kernel_path_surface.h"
#include "kernel_path_volume.h"
#include "kernel_path_subsurface.h"
#ifdef __KERNEL_DEBUG__
# include "kernel_debug.h"
@@ -413,172 +414,6 @@ ccl_device void kernel_path_indirect(KernelGlobals *kg,
}
}
#ifdef __SUBSURFACE__
# ifndef __KERNEL_CUDA__
ccl_device
# else
ccl_device_inline
# endif
bool kernel_path_subsurface_scatter(
KernelGlobals *kg,
ShaderData *sd,
ShaderData *emission_sd,
PathRadiance *L,
PathState *state,
RNG *rng,
Ray *ray,
float3 *throughput,
SubsurfaceIndirectRays *ss_indirect)
{
float bssrdf_probability;
ShaderClosure *sc = subsurface_scatter_pick_closure(kg, sd, &bssrdf_probability);
/* modify throughput for picking bssrdf or bsdf */
*throughput *= bssrdf_probability;
/* do bssrdf scatter step if we picked a bssrdf closure */
if(sc) {
/* We should never have two consecutive BSSRDF bounces,
* the second one should be converted to a diffuse BSDF to
* avoid this.
*/
kernel_assert(!ss_indirect->tracing);
uint lcg_state = lcg_state_init(rng, state, 0x68bc21eb);
SubsurfaceIntersection ss_isect;
float bssrdf_u, bssrdf_v;
path_state_rng_2D(kg, rng, state, PRNG_BSDF_U, &bssrdf_u, &bssrdf_v);
int num_hits = subsurface_scatter_multi_intersect(kg,
&ss_isect,
sd,
sc,
&lcg_state,
bssrdf_u, bssrdf_v,
false);
# ifdef __VOLUME__
ss_indirect->need_update_volume_stack =
kernel_data.integrator.use_volumes &&
sd->object_flag & SD_OBJECT_INTERSECTS_VOLUME;
# endif /* __VOLUME__ */
/* compute lighting with the BSDF closure */
for(int hit = 0; hit < num_hits; hit++) {
/* NOTE: We reuse the existing ShaderData, we assume the path
* integration loop stops when this function returns true.
*/
subsurface_scatter_multi_setup(kg,
&ss_isect,
hit,
sd,
state,
state->flag,
sc,
false);
PathState *hit_state = &ss_indirect->state[ss_indirect->num_rays];
Ray *hit_ray = &ss_indirect->rays[ss_indirect->num_rays];
float3 *hit_tp = &ss_indirect->throughputs[ss_indirect->num_rays];
PathRadiance *hit_L = &ss_indirect->L[ss_indirect->num_rays];
*hit_state = *state;
*hit_ray = *ray;
*hit_tp = *throughput;
hit_state->rng_offset += PRNG_BOUNCE_NUM;
path_radiance_init(hit_L, kernel_data.film.use_light_pass);
hit_L->direct_throughput = L->direct_throughput;
path_radiance_copy_indirect(hit_L, L);
kernel_path_surface_connect_light(kg, rng, sd, emission_sd, *hit_tp, state, hit_L);
if(kernel_path_surface_bounce(kg,
rng,
sd,
hit_tp,
hit_state,
hit_L,
hit_ray))
{
# ifdef __LAMP_MIS__
hit_state->ray_t = 0.0f;
# endif /* __LAMP_MIS__ */
# ifdef __VOLUME__
if(ss_indirect->need_update_volume_stack) {
Ray volume_ray = *ray;
/* Setup ray from previous surface point to the new one. */
volume_ray.D = normalize_len(hit_ray->P - volume_ray.P,
&volume_ray.t);
kernel_volume_stack_update_for_subsurface(
kg,
emission_sd,
&volume_ray,
hit_state->volume_stack);
}
# endif /* __VOLUME__ */
path_radiance_reset_indirect(L);
ss_indirect->num_rays++;
}
else {
path_radiance_accum_sample(L, hit_L, 1);
}
}
return true;
}
return false;
}
ccl_device_inline void kernel_path_subsurface_init_indirect(
SubsurfaceIndirectRays *ss_indirect)
{
ss_indirect->tracing = false;
ss_indirect->num_rays = 0;
}
ccl_device void kernel_path_subsurface_accum_indirect(
SubsurfaceIndirectRays *ss_indirect,
PathRadiance *L)
{
if(ss_indirect->tracing) {
path_radiance_sum_indirect(L);
path_radiance_accum_sample(&ss_indirect->direct_L, L, 1);
if(ss_indirect->num_rays == 0) {
*L = ss_indirect->direct_L;
}
}
}
ccl_device void kernel_path_subsurface_setup_indirect(
KernelGlobals *kg,
SubsurfaceIndirectRays *ss_indirect,
PathState *state,
Ray *ray,
PathRadiance *L,
float3 *throughput)
{
if(!ss_indirect->tracing) {
ss_indirect->direct_L = *L;
}
ss_indirect->tracing = true;
/* Setup state, ray and throughput for indirect SSS rays. */
ss_indirect->num_rays--;
Ray *indirect_ray = &ss_indirect->rays[ss_indirect->num_rays];
PathRadiance *indirect_L = &ss_indirect->L[ss_indirect->num_rays];
*state = ss_indirect->state[ss_indirect->num_rays];
*ray = *indirect_ray;
*L = *indirect_L;
*throughput = ss_indirect->throughputs[ss_indirect->num_rays];
state->rng_offset += ss_indirect->num_rays * PRNG_BOUNCE_NUM;
}
#endif /* __SUBSURFACE__ */
ccl_device_inline float4 kernel_path_integrate(KernelGlobals *kg,
RNG *rng,

187
intern/cycles/kernel/kernel_path_subsurface.h Normal file
View File

@@ -0,0 +1,187 @@
/*
* Copyright 2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
#ifdef __SUBSURFACE__
# ifndef __KERNEL_CUDA__
ccl_device
# else
ccl_device_inline
# endif
bool kernel_path_subsurface_scatter(
KernelGlobals *kg,
ShaderData *sd,
ShaderData *emission_sd,
PathRadiance *L,
ccl_addr_space PathState *state,
ccl_addr_space RNG *rng,
ccl_addr_space Ray *ray,
ccl_addr_space float3 *throughput,
ccl_addr_space SubsurfaceIndirectRays *ss_indirect)
{
float bssrdf_probability;
ShaderClosure *sc = subsurface_scatter_pick_closure(kg, sd, &bssrdf_probability);
/* modify throughput for picking bssrdf or bsdf */
*throughput *= bssrdf_probability;
/* do bssrdf scatter step if we picked a bssrdf closure */
if(sc) {
/* We should never have two consecutive BSSRDF bounces,
* the second one should be converted to a diffuse BSDF to
* avoid this.
*/
kernel_assert(!ss_indirect->tracing);
uint lcg_state = lcg_state_init_addrspace(rng, state, 0x68bc21eb);
SubsurfaceIntersection ss_isect;
float bssrdf_u, bssrdf_v;
path_state_rng_2D(kg, rng, state, PRNG_BSDF_U, &bssrdf_u, &bssrdf_v);
int num_hits = subsurface_scatter_multi_intersect(kg,
&ss_isect,
sd,
sc,
&lcg_state,
bssrdf_u, bssrdf_v,
false);
# ifdef __VOLUME__
ss_indirect->need_update_volume_stack =
kernel_data.integrator.use_volumes &&
sd->object_flag & SD_OBJECT_INTERSECTS_VOLUME;
# endif /* __VOLUME__ */
/* compute lighting with the BSDF closure */
for(int hit = 0; hit < num_hits; hit++) {
/* NOTE: We reuse the existing ShaderData, we assume the path
* integration loop stops when this function returns true.
*/
subsurface_scatter_multi_setup(kg,
&ss_isect,
hit,
sd,
state,
state->flag,
sc,
false);
ccl_addr_space PathState *hit_state = &ss_indirect->state[ss_indirect->num_rays];
ccl_addr_space Ray *hit_ray = &ss_indirect->rays[ss_indirect->num_rays];
ccl_addr_space float3 *hit_tp = &ss_indirect->throughputs[ss_indirect->num_rays];
PathRadiance *hit_L = &ss_indirect->L[ss_indirect->num_rays];
*hit_state = *state;
*hit_ray = *ray;
*hit_tp = *throughput;
hit_state->rng_offset += PRNG_BOUNCE_NUM;
path_radiance_init(hit_L, kernel_data.film.use_light_pass);
hit_L->direct_throughput = L->direct_throughput;
path_radiance_copy_indirect(hit_L, L);
kernel_path_surface_connect_light(kg, rng, sd, emission_sd, *hit_tp, state, hit_L);
if(kernel_path_surface_bounce(kg,
rng,
sd,
hit_tp,
hit_state,
hit_L,
hit_ray))
{
# ifdef __LAMP_MIS__
hit_state->ray_t = 0.0f;
# endif /* __LAMP_MIS__ */
# ifdef __VOLUME__
if(ss_indirect->need_update_volume_stack) {
Ray volume_ray = *ray;
/* Setup ray from previous surface point to the new one. */
volume_ray.D = normalize_len(hit_ray->P - volume_ray.P,
&volume_ray.t);
kernel_volume_stack_update_for_subsurface(
kg,
emission_sd,
&volume_ray,
hit_state->volume_stack);
}
# endif /* __VOLUME__ */
path_radiance_reset_indirect(L);
ss_indirect->num_rays++;
}
else {
path_radiance_accum_sample(L, hit_L, 1);
}
}
return true;
}
return false;
}
ccl_device_inline void kernel_path_subsurface_init_indirect(
ccl_addr_space SubsurfaceIndirectRays *ss_indirect)
{
ss_indirect->tracing = false;
ss_indirect->num_rays = 0;
}
ccl_device void kernel_path_subsurface_accum_indirect(
ccl_addr_space SubsurfaceIndirectRays *ss_indirect,
PathRadiance *L)
{
if(ss_indirect->tracing) {
path_radiance_sum_indirect(L);
path_radiance_accum_sample(&ss_indirect->direct_L, L, 1);
if(ss_indirect->num_rays == 0) {
*L = ss_indirect->direct_L;
}
}
}
ccl_device void kernel_path_subsurface_setup_indirect(
KernelGlobals *kg,
ccl_addr_space SubsurfaceIndirectRays *ss_indirect,
ccl_addr_space PathState *state,
ccl_addr_space Ray *ray,
PathRadiance *L,
ccl_addr_space float3 *throughput)
{
if(!ss_indirect->tracing) {
ss_indirect->direct_L = *L;
}
ss_indirect->tracing = true;
/* Setup state, ray and throughput for indirect SSS rays. */
ss_indirect->num_rays--;
ccl_addr_space Ray *indirect_ray = &ss_indirect->rays[ss_indirect->num_rays];
PathRadiance *indirect_L = &ss_indirect->L[ss_indirect->num_rays];
*state = ss_indirect->state[ss_indirect->num_rays];
*ray = *indirect_ray;
*L = *indirect_L;
*throughput = ss_indirect->throughputs[ss_indirect->num_rays];
state->rng_offset += ss_indirect->num_rays * PRNG_BOUNCE_NUM;
}
#endif /* __SUBSURFACE__ */
CCL_NAMESPACE_END

4
intern/cycles/kernel/kernel_path_surface.h
View File

@@ -16,7 +16,7 @@
CCL_NAMESPACE_BEGIN
#if defined(__BRANCHED_PATH__) || defined(__SUBSURFACE__)
#if (defined(__BRANCHED_PATH__) || defined(__SUBSURFACE__)) && !defined(__SPLIT_KERNEL__)
/* branched path tracing: connect path directly to position on one or more lights and add it to L */
ccl_device_noinline void kernel_branched_path_surface_connect_light(KernelGlobals *kg, RNG *rng,
@@ -188,7 +188,6 @@ ccl_device bool kernel_branched_path_surface_bounce(KernelGlobals *kg, RNG *rng,
#endif
#ifndef __SPLIT_KERNEL__
/* path tracing: connect path directly to position on a light and add it to L */
ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg, ccl_addr_space RNG *rng,
ShaderData *sd, ShaderData *emission_sd, float3 throughput, ccl_addr_space PathState *state,
@@ -226,7 +225,6 @@ ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg, ccl_
}
#endif
}
#endif
/* path tracing: bounce off or through surface to with new direction stored in ray */
ccl_device bool kernel_path_surface_bounce(KernelGlobals *kg,

22
intern/cycles/kernel/kernel_path_volume.h
View File

@@ -20,11 +20,11 @@ CCL_NAMESPACE_BEGIN
ccl_device_inline void kernel_path_volume_connect_light(
KernelGlobals *kg,
RNG *rng,
ccl_addr_space RNG *rng,
ShaderData *sd,
ShaderData *emission_sd,
float3 throughput,
PathState *state,
ccl_addr_space PathState *state,
PathRadiance *L)
{
#ifdef __EMISSION__
@@ -59,7 +59,7 @@ ccl_device_inline void kernel_path_volume_connect_light(
}
}
}
#endif
#endif /* __EMISSION__ */
}
#ifdef __KERNEL_GPU__
@@ -67,8 +67,14 @@ ccl_device_noinline
#else
ccl_device
#endif
bool kernel_path_volume_bounce(KernelGlobals *kg, RNG *rng,
ShaderData *sd, float3 *throughput, PathState *state, PathRadiance *L, Ray *ray)
bool kernel_path_volume_bounce(
KernelGlobals *kg,
ccl_addr_space RNG *rng,
ShaderData *sd,
ccl_addr_space float3 *throughput,
ccl_addr_space PathState *state,
PathRadiance *L,
ccl_addr_space Ray *ray)
{
/* sample phase function */
float phase_pdf;
@@ -111,6 +117,7 @@ bool kernel_path_volume_bounce(KernelGlobals *kg, RNG *rng,
return true;
}
#ifdef __BRANCHED_PATH__
ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg, RNG *rng,
ShaderData *sd, ShaderData *emission_sd, float3 throughput, PathState *state, PathRadiance *L,
bool sample_all_lights, Ray *ray, const VolumeSegment *segment)
@@ -261,10 +268,11 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg, RNG
}
}
}
#endif
#endif /* __EMISSION__ */
}
#endif /* __BRANCHED_PATH__ */
#endif
#endif /* __VOLUME_SCATTER__ */
CCL_NAMESPACE_END

14
intern/cycles/kernel/kernel_shader.h
View File

@@ -203,11 +203,11 @@ void shader_setup_from_subsurface(
# ifdef __INSTANCING__
if(isect->object != OBJECT_NONE) {
/* instance transform */
object_normal_transform(kg, sd, &sd->N);
object_normal_transform(kg, sd, &sd->Ng);
object_normal_transform_auto(kg, sd, &sd->N);
object_normal_transform_auto(kg, sd, &sd->Ng);
# ifdef __DPDU__
object_dir_transform(kg, sd, &sd->dPdu);
object_dir_transform(kg, sd, &sd->dPdv);
object_dir_transform_auto(kg, sd, &sd->dPdu);
object_dir_transform_auto(kg, sd, &sd->dPdv);
# endif
}
# endif
@@ -816,7 +816,7 @@ ccl_device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_b
*N_ = (is_zero(N))? sd->N: normalize(N);
if(texture_blur_)
*texture_blur_ = texture_blur/weight_sum;
*texture_blur_ = safe_divide(texture_blur, weight_sum);
return eval;
}
@@ -1036,8 +1036,8 @@ ccl_device int shader_phase_sample_closure(KernelGlobals *kg, const ShaderData *
ccl_device_inline void shader_eval_volume(KernelGlobals *kg,
ShaderData *sd,
PathState *state,
VolumeStack *stack,
ccl_addr_space PathState *state,
ccl_addr_space VolumeStack *stack,
int path_flag,
ShaderContext ctx)
{

59
intern/cycles/kernel/kernel_shadow.h
View File

@@ -24,7 +24,7 @@ ccl_device_forceinline bool shadow_handle_transparent_isect(
ShaderData *shadow_sd,
ccl_addr_space PathState *state,
# ifdef __VOLUME__
struct PathState *volume_state,
ccl_addr_space struct PathState *volume_state,
# endif
Intersection *isect,
Ray *ray,
@@ -152,7 +152,13 @@ ccl_device bool shadow_blocked_transparent_all_loop(KernelGlobals *kg,
int bounce = state->transparent_bounce;
Intersection *isect = hits;
# ifdef __VOLUME__
PathState ps = *state;
# ifdef __SPLIT_KERNEL__
ccl_addr_space PathState *ps = &kernel_split_state.state_shadow[ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0)];
# else
PathState ps_object;
PathState *ps = &ps_object;
# endif
*ps = *state;
# endif
sort_intersections(hits, num_hits);
for(int hit = 0; hit < num_hits; hit++, isect++) {
@@ -171,7 +177,7 @@ ccl_device bool shadow_blocked_transparent_all_loop(KernelGlobals *kg,
shadow_sd,
state,
#ifdef __VOLUME__
&ps,
ps,
#endif
isect,
ray,
@@ -188,8 +194,8 @@ ccl_device bool shadow_blocked_transparent_all_loop(KernelGlobals *kg,
}
# ifdef __VOLUME__
/* Attenuation for last line segment towards light. */
if(ps.volume_stack[0].shader != SHADER_NONE) {
kernel_volume_shadow(kg, shadow_sd, &ps, ray, &throughput);
if(ps->volume_stack[0].shader != SHADER_NONE) {
kernel_volume_shadow(kg, shadow_sd, ps, ray, &throughput);
}
# endif
*shadow = throughput;
@@ -214,7 +220,10 @@ ccl_device bool shadow_blocked_transparent_all(KernelGlobals *kg,
uint max_hits,
float3 *shadow)
{
# ifdef __KERNEL_CUDA__
# ifdef __SPLIT_KERNEL__
Intersection hits_[SHADOW_STACK_MAX_HITS];
Intersection *hits = &hits_[0];
# elif defined(__KERNEL_CUDA__)
Intersection *hits = kg->hits_stack;
# else
Intersection hits_stack[SHADOW_STACK_MAX_HITS];
@@ -276,7 +285,13 @@ ccl_device bool shadow_blocked_transparent_stepped_loop(
float3 Pend = ray->P + ray->D*ray->t;
int bounce = state->transparent_bounce;
# ifdef __VOLUME__
PathState ps = *state;
# ifdef __SPLIT_KERNEL__
ccl_addr_space PathState *ps = &kernel_split_state.state_shadow[ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0)];
# else
PathState ps_object;
PathState *ps = &ps_object;
# endif
*ps = *state;
# endif
for(;;) {
if(bounce >= kernel_data.integrator.transparent_max_bounce) {
@@ -299,7 +314,7 @@ ccl_device bool shadow_blocked_transparent_stepped_loop(
shadow_sd,
state,
#ifdef __VOLUME__
&ps,
ps,
#endif
isect,
ray,
@@ -316,8 +331,8 @@ ccl_device bool shadow_blocked_transparent_stepped_loop(
}
# ifdef __VOLUME__
/* Attenuation for last line segment towards light. */
if(ps.volume_stack[0].shader != SHADER_NONE) {
kernel_volume_shadow(kg, shadow_sd, &ps, ray, &throughput);
if(ps->volume_stack[0].shader != SHADER_NONE) {
kernel_volume_shadow(kg, shadow_sd, ps, ray, &throughput);
}
# endif
*shadow *= throughput;
@@ -365,21 +380,11 @@ ccl_device bool shadow_blocked_transparent_stepped(
ccl_device_inline bool shadow_blocked(KernelGlobals *kg,
ShaderData *shadow_sd,
ccl_addr_space PathState *state,
ccl_addr_space Ray *ray_input,
Ray *ray_input,
float3 *shadow)
{
/* Special trickery for split kernel: some data is coming from the
* global memory.
*/
#ifdef __SPLIT_KERNEL__
Ray private_ray = *ray_input;
Ray *ray = &private_ray;
Intersection *isect = &kernel_split_state.isect_shadow[ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0)];
#else /* __SPLIT_KERNEL__ */
Ray *ray = ray_input;
Intersection isect_object;
Intersection *isect = &isect_object;
#endif /* __SPLIT_KERNEL__ */
Intersection isect;
/* Some common early checks. */
*shadow = make_float3(1.0f, 1.0f, 1.0f);
if(ray->t == 0.0f) {
@@ -397,7 +402,7 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg,
shadow_sd,
state,
ray,
isect,
&isect,
shadow);
}
#ifdef __TRANSPARENT_SHADOWS__
@@ -423,11 +428,11 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg,
const bool blocked = scene_intersect(kg,
*ray,
PATH_RAY_SHADOW_OPAQUE,
isect,
&isect,
NULL,
0.0f, 0.0f);
const bool is_transparent_isect = blocked
? shader_transparent_shadow(kg, isect)
? shader_transparent_shadow(kg, &isect)
: false;
if(!blocked || !is_transparent_isect ||
max_hits + 1 >= SHADOW_STACK_MAX_HITS)
@@ -436,7 +441,7 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg,
shadow_sd,
state,
ray,
isect,
&isect,
blocked,
is_transparent_isect,
shadow);
@@ -454,7 +459,7 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg,
shadow_sd,
state,
ray,
isect,
&isect,
shadow);
# endif /* __SHADOW_RECORD_ALL__ */
#endif /* __TRANSPARENT_SHADOWS__ */

23
intern/cycles/kernel/kernel_subsurface.h
View File

@@ -185,7 +185,7 @@ ccl_device float3 subsurface_color_pow(float3 color, float exponent)
ccl_device void subsurface_color_bump_blur(KernelGlobals *kg,
ShaderData *sd,
PathState *state,
ccl_addr_space PathState *state,
int state_flag,
float3 *eval,
float3 *N)
@@ -277,7 +277,12 @@ ccl_device_inline int subsurface_scatter_multi_intersect(
float3 disk_P = (disk_r*cosf(phi)) * disk_T + (disk_r*sinf(phi)) * disk_B;
/* create ray */
#ifdef __SPLIT_KERNEL__
Ray ray_object = ss_isect->ray;
Ray *ray = &ray_object;
#else
Ray *ray = &ss_isect->ray;
#endif
ray->P = sd->P + disk_N*disk_height + disk_P;
ray->D = -disk_N;
ray->t = 2.0f*disk_height;
@@ -351,6 +356,10 @@ ccl_device_inline int subsurface_scatter_multi_intersect(
ss_isect->weight[hit] = eval;
}
#ifdef __SPLIT_KERNEL__
ss_isect->ray = *ray;
#endif
return num_eval_hits;
}
@@ -359,13 +368,19 @@ ccl_device_noinline void subsurface_scatter_multi_setup(
SubsurfaceIntersection* ss_isect,
int hit,
ShaderData *sd,
PathState *state,
ccl_addr_space PathState *state,
int state_flag,
ShaderClosure *sc,
bool all)
{
#ifdef __SPLIT_KERNEL__
Ray ray_object = ss_isect->ray;
Ray *ray = &ray_object;
#else
Ray *ray = &ss_isect->ray;
#endif
/* Setup new shading point. */
shader_setup_from_subsurface(kg, sd, &ss_isect->hits[hit], &ss_isect->ray);
shader_setup_from_subsurface(kg, sd, &ss_isect->hits[hit], ray);
/* Optionally blur colors and bump mapping. */
float3 weight = ss_isect->weight[hit];
@@ -376,6 +391,7 @@ ccl_device_noinline void subsurface_scatter_multi_setup(
subsurface_scatter_setup_diffuse_bsdf(sd, weight, true, N);
}
#ifndef __SPLIT_KERNEL__
/* subsurface scattering step, from a point on the surface to another nearby point on the same object */
ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd, PathState *state,
int state_flag, ShaderClosure *sc, uint *lcg_state, float disk_u, float disk_v, bool all)
@@ -465,6 +481,7 @@ ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd, PathS
/* setup diffuse bsdf */
subsurface_scatter_setup_diffuse_bsdf(sd, eval, (ss_isect.num_hits > 0), N);
}
#endif /* ! __SPLIT_KERNEL__ */
CCL_NAMESPACE_END

24
intern/cycles/kernel/kernel_types.h
View File

@@ -76,15 +76,13 @@ CCL_NAMESPACE_BEGIN
# ifdef WITH_OSL
# define __OSL__
# endif
# ifndef __SPLIT_KERNEL__
# define __SUBSURFACE__
# endif
# define __SUBSURFACE__
# define __CMJ__
# define __VOLUME__
# define __VOLUME_SCATTER__
# define __SHADOW_RECORD_ALL__
# ifndef __SPLIT_KERNEL__
# define __VOLUME__
# define __VOLUME_DECOUPLED__
# define __VOLUME_SCATTER__
# define __SHADOW_RECORD_ALL__
# define __VOLUME_RECORD_ALL__
# endif
#endif /* __KERNEL_CPU__ */
@@ -130,6 +128,10 @@ CCL_NAMESPACE_BEGIN
# define __CL_USE_NATIVE__
# define __KERNEL_SHADING__
# define __KERNEL_ADV_SHADING__
# define __SUBSURFACE__
# define __VOLUME__
# define __VOLUME_SCATTER__
# define __SHADOW_RECORD_ALL__
# endif /* __KERNEL_OPENCL_AMD__ */
# ifdef __KERNEL_OPENCL_INTEL_CPU__
@@ -552,7 +554,7 @@ typedef struct Ray {
/* Intersection */
typedef ccl_addr_space struct Intersection {
typedef struct Intersection {
float t, u, v;
int prim;
int object;
@@ -934,7 +936,7 @@ typedef struct PathState {
/* Subsurface */
/* Struct to gather multiple SSS hits. */
struct SubsurfaceIntersection
typedef struct SubsurfaceIntersection
{
Ray ray;
float3 weight[BSSRDF_MAX_HITS];
@@ -942,10 +944,10 @@ struct SubsurfaceIntersection
int num_hits;
struct Intersection hits[BSSRDF_MAX_HITS];
float3 Ng[BSSRDF_MAX_HITS];
};
} SubsurfaceIntersection;
/* Struct to gather SSS indirect rays and delay tracing them. */
struct SubsurfaceIndirectRays
typedef struct SubsurfaceIndirectRays
{
bool need_update_volume_stack;
bool tracing;
@@ -956,7 +958,7 @@ struct SubsurfaceIndirectRays
struct Ray rays[BSSRDF_MAX_HITS];
float3 throughputs[BSSRDF_MAX_HITS];
struct PathRadiance L[BSSRDF_MAX_HITS];
};
} SubsurfaceIndirectRays;
/* Constant Kernel Data
*

79
intern/cycles/kernel/kernel_volume.h
View File

@@ -38,7 +38,7 @@ typedef struct VolumeShaderCoefficients {
/* evaluate shader to get extinction coefficient at P */
ccl_device_inline bool volume_shader_extinction_sample(KernelGlobals *kg,
ShaderData *sd,
PathState *state,
ccl_addr_space PathState *state,
float3 P,
float3 *extinction)
{
@@ -64,7 +64,7 @@ ccl_device_inline bool volume_shader_extinction_sample(KernelGlobals *kg,
/* evaluate shader to get absorption, scattering and emission at P */
ccl_device_inline bool volume_shader_sample(KernelGlobals *kg,
ShaderData *sd,
PathState *state,
ccl_addr_space PathState *state,
float3 P,
VolumeShaderCoefficients *coeff)
{
@@ -112,7 +112,7 @@ ccl_device float kernel_volume_channel_get(float3 value, int channel)
return (channel == 0)? value.x: ((channel == 1)? value.y: value.z);
}
ccl_device bool volume_stack_is_heterogeneous(KernelGlobals *kg, VolumeStack *stack)
ccl_device bool volume_stack_is_heterogeneous(KernelGlobals *kg, ccl_addr_space VolumeStack *stack)
{
for(int i = 0; stack[i].shader != SHADER_NONE; i++) {
int shader_flag = kernel_tex_fetch(__shader_flag, (stack[i].shader & SHADER_MASK)*SHADER_SIZE);
@@ -161,7 +161,11 @@ ccl_device int volume_stack_sampling_method(KernelGlobals *kg, VolumeStack *stac
/* homogeneous volume: assume shader evaluation at the starts gives
* the extinction coefficient for the entire line segment */
ccl_device void kernel_volume_shadow_homogeneous(KernelGlobals *kg, PathState *state, Ray *ray, ShaderData *sd, float3 *throughput)
ccl_device void kernel_volume_shadow_homogeneous(KernelGlobals *kg,
ccl_addr_space PathState *state,
Ray *ray,
ShaderData *sd,
float3 *throughput)
{
float3 sigma_t;
@@ -171,7 +175,11 @@ ccl_device void kernel_volume_shadow_homogeneous(KernelGlobals *kg, PathState *s
/* heterogeneous volume: integrate stepping through the volume until we
* reach the end, get absorbed entirely, or run out of iterations */
ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg, PathState *state, Ray *ray, ShaderData *sd, float3 *throughput)
ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg,
ccl_addr_space PathState *state,
Ray *ray,
ShaderData *sd,
float3 *throughput)
{
float3 tp = *throughput;
const float tp_eps = 1e-6f; /* todo: this is likely not the right value */
@@ -179,7 +187,7 @@ ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg, PathState
/* prepare for stepping */
int max_steps = kernel_data.integrator.volume_max_steps;
float step = kernel_data.integrator.volume_step_size;
float random_jitter_offset = lcg_step_float(&state->rng_congruential) * step;
float random_jitter_offset = lcg_step_float_addrspace(&state->rng_congruential) * step;
/* compute extinction at the start */
float t = 0.0f;
@@ -193,7 +201,7 @@ ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg, PathState
/* use random position inside this segment to sample shader */
if(new_t == ray->t)
random_jitter_offset = lcg_step_float(&state->rng_congruential) * dt;
random_jitter_offset = lcg_step_float_addrspace(&state->rng_congruential) * dt;
float3 new_P = ray->P + ray->D * (t + random_jitter_offset);
float3 sigma_t;
@@ -227,7 +235,11 @@ ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg, PathState
/* get the volume attenuation over line segment defined by ray, with the
* assumption that there are no surfaces blocking light between the endpoints */
ccl_device_noinline void kernel_volume_shadow(KernelGlobals *kg, ShaderData *shadow_sd, PathState *state, Ray *ray, float3 *throughput)
ccl_device_noinline void kernel_volume_shadow(KernelGlobals *kg,
ShaderData *shadow_sd,
ccl_addr_space PathState *state,
Ray *ray,
float3 *throughput)
{
shader_setup_from_volume(kg, shadow_sd, ray);
@@ -341,9 +353,15 @@ ccl_device float3 kernel_volume_emission_integrate(VolumeShaderCoefficients *coe
/* homogeneous volume: assume shader evaluation at the start gives
* the volume shading coefficient for the entire line segment */
ccl_device VolumeIntegrateResult kernel_volume_integrate_homogeneous(KernelGlobals *kg,
PathState *state, Ray *ray, ShaderData *sd, PathRadiance *L, float3 *throughput,
RNG *rng, bool probalistic_scatter)
ccl_device VolumeIntegrateResult kernel_volume_integrate_homogeneous(
KernelGlobals *kg,
ccl_addr_space PathState *state,
Ray *ray,
ShaderData *sd,
PathRadiance *L,
ccl_addr_space float3 *throughput,
ccl_addr_space RNG *rng,
bool probalistic_scatter)
{
VolumeShaderCoefficients coeff;
@@ -444,8 +462,14 @@ ccl_device VolumeIntegrateResult kernel_volume_integrate_homogeneous(KernelGloba
* volume until we reach the end, get absorbed entirely, or run out of
* iterations. this does probabilistically scatter or get transmitted through
* for path tracing where we don't want to branch. */
ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(KernelGlobals *kg,
PathState *state, Ray *ray, ShaderData *sd, PathRadiance *L, float3 *throughput, RNG *rng)
ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(
KernelGlobals *kg,
ccl_addr_space PathState *state,
Ray *ray,
ShaderData *sd,
PathRadiance *L,
ccl_addr_space float3 *throughput,
ccl_addr_space RNG *rng)
{
float3 tp = *throughput;
const float tp_eps = 1e-6f; /* todo: this is likely not the right value */
@@ -453,7 +477,7 @@ ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(
/* prepare for stepping */
int max_steps = kernel_data.integrator.volume_max_steps;
float step_size = kernel_data.integrator.volume_step_size;
float random_jitter_offset = lcg_step_float(&state->rng_congruential) * step_size;
float random_jitter_offset = lcg_step_float_addrspace(&state->rng_congruential) * step_size;
/* compute coefficients at the start */
float t = 0.0f;
@@ -474,7 +498,7 @@ ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(
/* use random position inside this segment to sample shader */
if(new_t == ray->t)
random_jitter_offset = lcg_step_float(&state->rng_congruential) * dt;
random_jitter_offset = lcg_step_float_addrspace(&state->rng_congruential) * dt;
float3 new_P = ray->P + ray->D * (t + random_jitter_offset);
VolumeShaderCoefficients coeff;
@@ -579,8 +603,15 @@ ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(
* ray, with the assumption that there are no surfaces blocking light
* between the endpoints. distance sampling is used to decide if we will
* scatter or not. */
ccl_device_noinline VolumeIntegrateResult kernel_volume_integrate(KernelGlobals *kg,
PathState *state, ShaderData *sd, Ray *ray, PathRadiance *L, float3 *throughput, RNG *rng, bool heterogeneous)
ccl_device_noinline VolumeIntegrateResult kernel_volume_integrate(
KernelGlobals *kg,
ccl_addr_space PathState *state,
ShaderData *sd,
Ray *ray,
PathRadiance *L,
ccl_addr_space float3 *throughput,
ccl_addr_space RNG *rng,
bool heterogeneous)
{
shader_setup_from_volume(kg, sd, ray);
@@ -590,6 +621,7 @@ ccl_device_noinline VolumeIntegrateResult kernel_volume_integrate(KernelGlobals
return kernel_volume_integrate_homogeneous(kg, state, ray, sd, L, throughput, rng, true);
}
#ifndef __SPLIT_KERNEL__
/* Decoupled Volume Sampling
*
* VolumeSegment is list of coefficients and transmittance stored at all steps
@@ -990,6 +1022,7 @@ ccl_device VolumeIntegrateResult kernel_volume_decoupled_scatter(
return VOLUME_PATH_SCATTERED;
}
#endif /* __SPLIT_KERNEL */
/* decide if we need to use decoupled or not */
ccl_device bool kernel_volume_use_decoupled(KernelGlobals *kg, bool heterogeneous, bool direct, int sampling_method)
@@ -1021,9 +1054,9 @@ ccl_device bool kernel_volume_use_decoupled(KernelGlobals *kg, bool heterogeneou
ccl_device void kernel_volume_stack_init(KernelGlobals *kg,
ShaderData *stack_sd,
const PathState *state,
const Ray *ray,
VolumeStack *stack)
ccl_addr_space const PathState *state,
ccl_addr_space const Ray *ray,
ccl_addr_space VolumeStack *stack)
{
/* NULL ray happens in the baker, does it need proper initialization of
* camera in volume?
@@ -1166,7 +1199,7 @@ ccl_device void kernel_volume_stack_init(KernelGlobals *kg,
}
}
ccl_device void kernel_volume_stack_enter_exit(KernelGlobals *kg, ShaderData *sd, VolumeStack *stack)
ccl_device void kernel_volume_stack_enter_exit(KernelGlobals *kg, ShaderData *sd, ccl_addr_space VolumeStack *stack)
{
/* todo: we should have some way for objects to indicate if they want the
* world shader to work inside them. excluding it by default is problematic
@@ -1215,7 +1248,7 @@ ccl_device void kernel_volume_stack_enter_exit(KernelGlobals *kg, ShaderData *sd
ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg,
ShaderData *stack_sd,
Ray *ray,
VolumeStack *stack)
ccl_addr_space VolumeStack *stack)
{
kernel_assert(kernel_data.integrator.use_volumes);
@@ -1277,7 +1310,7 @@ ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg,
* the world's one after the last bounce to avoid render artifacts.
*/
ccl_device_inline void kernel_volume_clean_stack(KernelGlobals *kg,
VolumeStack *volume_stack)
ccl_addr_space VolumeStack *volume_stack)
{
if(kernel_data.background.volume_shader != SHADER_NONE) {
/* Keep the world's volume in stack. */

9
intern/cycles/kernel/kernels/cpu/kernel_cpu.h
View File

@@ -74,13 +74,18 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
DECLARE_SPLIT_KERNEL_FUNCTION(path_init)
DECLARE_SPLIT_KERNEL_FUNCTION(scene_intersect)
DECLARE_SPLIT_KERNEL_FUNCTION(lamp_emission)
DECLARE_SPLIT_KERNEL_FUNCTION(do_volume)
DECLARE_SPLIT_KERNEL_FUNCTION(queue_enqueue)
DECLARE_SPLIT_KERNEL_FUNCTION(background_buffer_update)
DECLARE_SPLIT_KERNEL_FUNCTION(indirect_background)
DECLARE_SPLIT_KERNEL_FUNCTION(shader_eval)
DECLARE_SPLIT_KERNEL_FUNCTION(holdout_emission_blurring_pathtermination_ao)
DECLARE_SPLIT_KERNEL_FUNCTION(subsurface_scatter)
DECLARE_SPLIT_KERNEL_FUNCTION(direct_lighting)
DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked)
DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao)
DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl)
DECLARE_SPLIT_KERNEL_FUNCTION(next_iteration_setup)
DECLARE_SPLIT_KERNEL_FUNCTION(indirect_subsurface)
DECLARE_SPLIT_KERNEL_FUNCTION(buffer_update)
void KERNEL_FUNCTION_FULL_NAME(register_functions)(void(*reg)(const char* name, void* func));

27
intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
View File

@@ -41,13 +41,18 @@
# include "split/kernel_path_init.h"
# include "split/kernel_scene_intersect.h"
# include "split/kernel_lamp_emission.h"
# include "split/kernel_do_volume.h"
# include "split/kernel_queue_enqueue.h"
# include "split/kernel_background_buffer_update.h"
# include "split/kernel_indirect_background.h"
# include "split/kernel_shader_eval.h"
# include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
# include "split/kernel_subsurface_scatter.h"
# include "split/kernel_direct_lighting.h"
# include "split/kernel_shadow_blocked.h"
# include "split/kernel_shadow_blocked_ao.h"
# include "split/kernel_shadow_blocked_dl.h"
# include "split/kernel_next_iteration_setup.h"
# include "split/kernel_indirect_subsurface.h"
# include "split/kernel_buffer_update.h"
#endif
CCL_NAMESPACE_BEGIN
@@ -166,13 +171,18 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
DEFINE_SPLIT_KERNEL_FUNCTION(path_init)
DEFINE_SPLIT_KERNEL_FUNCTION(scene_intersect)
DEFINE_SPLIT_KERNEL_FUNCTION(lamp_emission)
DEFINE_SPLIT_KERNEL_FUNCTION(do_volume)
DEFINE_SPLIT_KERNEL_FUNCTION(queue_enqueue)
DEFINE_SPLIT_KERNEL_FUNCTION(background_buffer_update)
DEFINE_SPLIT_KERNEL_FUNCTION(indirect_background)
DEFINE_SPLIT_KERNEL_FUNCTION(shader_eval)
DEFINE_SPLIT_KERNEL_FUNCTION(holdout_emission_blurring_pathtermination_ao)
DEFINE_SPLIT_KERNEL_FUNCTION(subsurface_scatter)
DEFINE_SPLIT_KERNEL_FUNCTION(direct_lighting)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl)
DEFINE_SPLIT_KERNEL_FUNCTION(next_iteration_setup)
DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface)
DEFINE_SPLIT_KERNEL_FUNCTION(buffer_update)
void KERNEL_FUNCTION_FULL_NAME(register_functions)(void(*reg)(const char* name, void* func))
{
@@ -189,13 +199,18 @@ void KERNEL_FUNCTION_FULL_NAME(register_functions)(void(*reg)(const char* name,
REGISTER(path_init);
REGISTER(scene_intersect);
REGISTER(lamp_emission);
REGISTER(do_volume);
REGISTER(queue_enqueue);
REGISTER(background_buffer_update);
REGISTER(indirect_background);
REGISTER(shader_eval);
REGISTER(holdout_emission_blurring_pathtermination_ao);
REGISTER(subsurface_scatter);
REGISTER(direct_lighting);
REGISTER(shadow_blocked);
REGISTER(shadow_blocked_ao);
REGISTER(shadow_blocked_dl);
REGISTER(next_iteration_setup);
REGISTER(indirect_subsurface);
REGISTER(buffer_update);
#undef REGISTER
#undef REGISTER_EVAL_NAME

6
intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl → intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
View File

@@ -16,11 +16,11 @@
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_shadow_blocked.h"
#include "split/kernel_buffer_update.h"
__kernel void kernel_ocl_path_trace_shadow_blocked(
__kernel void kernel_ocl_path_trace_buffer_update(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_shadow_blocked(kg);
kernel_buffer_update(kg);
}

26
intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl Normal file
View File

@@ -0,0 +1,26 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_do_volume.h"
__kernel void kernel_ocl_path_trace_do_volume(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_do_volume(kg);
}

26
intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl Normal file
View File

@@ -0,0 +1,26 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_indirect_background.h"
__kernel void kernel_ocl_path_trace_indirect_background(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_indirect_background(kg);
}

26
intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl Normal file
View File

@@ -0,0 +1,26 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_indirect_subsurface.h"
__kernel void kernel_ocl_path_trace_indirect_subsurface(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_indirect_subsurface(kg);
}

6
intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl → intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
View File

@@ -16,11 +16,11 @@
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_background_buffer_update.h"
#include "split/kernel_shadow_blocked_ao.h"
__kernel void kernel_ocl_path_trace_background_buffer_update(
__kernel void kernel_ocl_path_trace_shadow_blocked_ao(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_background_buffer_update(kg);
kernel_shadow_blocked_ao(kg);
}

26
intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl Normal file
View File

@@ -0,0 +1,26 @@
/*
* Copyright 2011-2015 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_shadow_blocked_dl.h"
__kernel void kernel_ocl_path_trace_shadow_blocked_dl(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_shadow_blocked_dl(kg);
}

35
intern/cycles/kernel/kernels/opencl/kernel_split.cl Normal file
View File

@@ -0,0 +1,35 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_state_buffer_size.cl"
#include "kernel_data_init.cl"
#include "kernel_path_init.cl"
#include "kernel_scene_intersect.cl"
#include "kernel_lamp_emission.cl"
#include "kernel_do_volume.cl"
#include "kernel_indirect_background.cl"
#include "kernel_queue_enqueue.cl"
#include "kernel_shader_eval.cl"
#include "kernel_holdout_emission_blurring_pathtermination_ao.cl"
#include "kernel_subsurface_scatter.cl"
#include "kernel_direct_lighting.cl"
#include "kernel_shadow_blocked_ao.cl"
#include "kernel_shadow_blocked_dl.cl"
#include "kernel_next_iteration_setup.cl"
#include "kernel_indirect_subsurface.cl"
#include "kernel_buffer_update.cl"

26
intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl Normal file
View File

@@ -0,0 +1,26 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "kernel_compat_opencl.h"
#include "split/kernel_split_common.h"
#include "split/kernel_subsurface_scatter.h"
__kernel void kernel_ocl_path_trace_subsurface_scatter(
KernelGlobals *kg,
ccl_constant KernelData *data)
{
kernel_subsurface_scatter(kg);
}

25
intern/cycles/kernel/split/kernel_background_buffer_update.h → intern/cycles/kernel/split/kernel_buffer_update.h
View File

@@ -69,7 +69,7 @@ CCL_NAMESPACE_BEGIN
* QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE and RAY_REGENERATED rays
* QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be empty
*/
ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
ccl_device void kernel_buffer_update(KernelGlobals *kg)
{
ccl_local unsigned int local_queue_atomics;
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
@@ -141,26 +141,6 @@ ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
rng_state += kernel_split_params.offset + pixel_x + pixel_y*stride;
buffer += (kernel_split_params.offset + pixel_x + pixel_y*stride) * kernel_data.film.pass_stride;
if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
/* eval background shader if nothing hit */
if(kernel_data.background.transparent && (state->flag & PATH_RAY_CAMERA)) {
*L_transparent = (*L_transparent) + average((*throughput));
#ifdef __PASSES__
if(!(kernel_data.film.pass_flag & PASS_BACKGROUND))
#endif
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
}
if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
#ifdef __BACKGROUND__
/* sample background shader */
float3 L_background = indirect_background(kg, &kernel_split_state.sd_DL_shadow[ray_index], state, ray);
path_radiance_accum_background(L, (*throughput), L_background, state->bounce);
#endif
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
}
}
if(IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) {
float3 L_sum = path_radiance_clamp_and_sum(kg, L);
kernel_write_light_passes(kg, buffer, L, sample);
@@ -207,6 +187,9 @@ ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
*L_transparent = 0.0f;
path_radiance_init(L, kernel_data.film.use_light_pass);
path_state_init(kg, &kernel_split_state.sd_DL_shadow[ray_index], state, rng, sample, ray);
#ifdef __SUBSURFACE__
kernel_path_subsurface_init_indirect(&kernel_split_state.ss_rays[ray_index]);
#endif
#ifdef __KERNEL_DEBUG__
debug_data_init(debug_data);
#endif

97
intern/cycles/kernel/split/kernel_do_volume.h Normal file
View File

@@ -0,0 +1,97 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
ccl_device void kernel_do_volume(KernelGlobals *kg)
{
#ifdef __VOLUME__
/* We will empty this queue in this kernel. */
if(ccl_global_id(0) == 0 && ccl_global_id(1) == 0) {
kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0;
}
/* Fetch use_queues_flag. */
ccl_local char local_use_queues_flag;
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
local_use_queues_flag = *kernel_split_params.use_queues_flag;
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
if(local_use_queues_flag) {
ray_index = get_ray_index(kg, ray_index,
QUEUE_ACTIVE_AND_REGENERATED_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
1);
if(ray_index == QUEUE_EMPTY_SLOT) {
return;
}
}
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE) ||
IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND)) {
bool hit = ! IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND);
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
ccl_global RNG *rng = &kernel_split_state.rng[ray_index];
ccl_global Intersection *isect = &kernel_split_state.isect[ray_index];
ShaderData *sd = &kernel_split_state.sd[ray_index];
ShaderData *sd_input = &kernel_split_state.sd_DL_shadow[ray_index];
/* Sanitize volume stack. */
if(!hit) {
kernel_volume_clean_stack(kg, state->volume_stack);
}
/* volume attenuation, emission, scatter */
if(state->volume_stack[0].shader != SHADER_NONE) {
Ray volume_ray = *ray;
volume_ray.t = (hit)? isect->t: FLT_MAX;
bool heterogeneous = volume_stack_is_heterogeneous(kg, state->volume_stack);
{
/* integrate along volume segment with distance sampling */
VolumeIntegrateResult result = kernel_volume_integrate(
kg, state, sd, &volume_ray, L, throughput, rng, heterogeneous);
# ifdef __VOLUME_SCATTER__
if(result == VOLUME_PATH_SCATTERED) {
/* direct lighting */
kernel_path_volume_connect_light(kg, rng, sd, sd_input, *throughput, state, L);
/* indirect light bounce */
if(kernel_path_volume_bounce(kg, rng, sd, throughput, state, L, ray))
ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED);
else
ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_UPDATE_BUFFER);
}
# endif
}
}
}
#endif
}
CCL_NAMESPACE_END

87
intern/cycles/kernel/split/kernel_indirect_background.h Normal file
View File

@@ -0,0 +1,87 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
ccl_device void kernel_indirect_background(KernelGlobals *kg)
{
/*
ccl_local unsigned int local_queue_atomics;
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
local_queue_atomics = 0;
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
// */
int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
ray_index = get_ray_index(kg, ray_index,
QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
0);
#ifdef __COMPUTE_DEVICE_GPU__
/* If we are executing on a GPU device, we exit all threads that are not
* required.
*
* If we are executing on a CPU device, then we need to keep all threads
* active since we have barrier() calls later in the kernel. CPU devices,
* expect all threads to execute barrier statement.
*/
if(ray_index == QUEUE_EMPTY_SLOT) {
return;
}
#endif
#ifndef __COMPUTE_DEVICE_GPU__
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
ccl_global char *ray_state = kernel_split_state.ray_state;
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
ccl_global float *L_transparent = &kernel_split_state.L_transparent[ray_index];
if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
/* eval background shader if nothing hit */
if(kernel_data.background.transparent && (state->flag & PATH_RAY_CAMERA)) {
*L_transparent = (*L_transparent) + average((*throughput));
#ifdef __PASSES__
if(!(kernel_data.film.pass_flag & PASS_BACKGROUND))
#endif
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
}
if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
#ifdef __BACKGROUND__
/* sample background shader */
float3 L_background = indirect_background(kg, &kernel_split_state.sd_DL_shadow[ray_index], state, ray);
path_radiance_accum_background(L, (*throughput), L_background, state->bounce);
#endif
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
}
}
#ifndef __COMPUTE_DEVICE_GPU__
}
#endif
}
CCL_NAMESPACE_END

77
intern/cycles/kernel/split/kernel_indirect_subsurface.h Normal file
View File

@@ -0,0 +1,77 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
ccl_device void kernel_indirect_subsurface(KernelGlobals *kg)
{
int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
if(thread_index == 0) {
/* We will empty both queues in this kernel. */
kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0;
kernel_split_params.queue_index[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = 0;
}
int ray_index;
get_ray_index(kg, thread_index,
QUEUE_ACTIVE_AND_REGENERATED_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
1);
ray_index = get_ray_index(kg, thread_index,
QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
1);
#ifdef __SUBSURFACE__
if(ray_index == QUEUE_EMPTY_SLOT) {
return;
}
ccl_global char *ray_state = kernel_split_state.ray_state;
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
if(IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) {
ccl_addr_space SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index];
kernel_path_subsurface_accum_indirect(ss_indirect, L);
/* Trace indirect subsurface rays by restarting the loop. this uses less
* stack memory than invoking kernel_path_indirect.
*/
if(ss_indirect->num_rays) {
kernel_path_subsurface_setup_indirect(kg,
ss_indirect,
state,
ray,
L,
throughput);
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED);
}
else {
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
}
}
#endif /* __SUBSURFACE__ */
}
CCL_NAMESPACE_END

9
intern/cycles/kernel/split/kernel_lamp_emission.h
View File

@@ -38,10 +38,12 @@ CCL_NAMESPACE_BEGIN
*/
ccl_device void kernel_lamp_emission(KernelGlobals *kg)
{
#ifndef __VOLUME__
/* We will empty this queue in this kernel. */
if(ccl_global_id(0) == 0 && ccl_global_id(1) == 0) {
kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0;
}
#endif
/* Fetch use_queues_flag. */
ccl_local char local_use_queues_flag;
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
@@ -55,7 +57,12 @@ ccl_device void kernel_lamp_emission(KernelGlobals *kg)
QUEUE_ACTIVE_AND_REGENERATED_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
1);
#ifndef __VOLUME__
1
#else
0
#endif
);
if(ray_index == QUEUE_EMPTY_SLOT) {
return;
}

4
intern/cycles/kernel/split/kernel_path_init.h
View File

@@ -82,6 +82,10 @@ ccl_device void kernel_path_init(KernelGlobals *kg) {
&kernel_split_state.rng[ray_index],
my_sample,
&kernel_split_state.ray[ray_index]);
#ifdef __SUBSURFACE__
kernel_path_subsurface_init_indirect(&kernel_split_state.ss_rays[ray_index]);
#endif
#ifdef __KERNEL_DEBUG__
debug_data_init(&kernel_split_state.debug_data[ray_index]);
#endif

6
intern/cycles/kernel/split/kernel_queue_enqueue.h
View File

@@ -63,10 +63,12 @@ ccl_device void kernel_queue_enqueue(KernelGlobals *kg)
int queue_number = -1;
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND)) {
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND) ||
IS_STATE(kernel_split_state.ray_state, ray_index, RAY_UPDATE_BUFFER)) {
queue_number = QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS;
}
else if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) {
else if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE) ||
IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) {
queue_number = QUEUE_ACTIVE_AND_REGENERATED_RAYS;
}

13
intern/cycles/kernel/split/kernel_scene_intersect.h
View File

@@ -93,7 +93,7 @@ ccl_device void kernel_scene_intersect(KernelGlobals *kg)
#ifdef __KERNEL_DEBUG__
DebugData *debug_data = &kernel_split_state.debug_data[ray_index];
#endif
Intersection *isect = &kernel_split_state.isect[ray_index];
Intersection isect;
PathState state = kernel_split_state.path_state[ray_index];
Ray ray = kernel_split_state.ray[ray_index];
@@ -116,16 +116,17 @@ ccl_device void kernel_scene_intersect(KernelGlobals *kg)
lcg_state = lcg_state_init(&rng, &state, 0x51633e2d);
}
bool hit = scene_intersect(kg, ray, visibility, isect, &lcg_state, difl, extmax);
bool hit = scene_intersect(kg, ray, visibility, &isect, &lcg_state, difl, extmax);
#else
bool hit = scene_intersect(kg, ray, visibility, isect, NULL, 0.0f, 0.0f);
bool hit = scene_intersect(kg, ray, visibility, &isect, NULL, 0.0f, 0.0f);
#endif
kernel_split_state.isect[ray_index] = isect;
#ifdef __KERNEL_DEBUG__
if(state.flag & PATH_RAY_CAMERA) {
debug_data->num_bvh_traversed_nodes += isect->num_traversed_nodes;
debug_data->num_bvh_traversed_instances += isect->num_traversed_instances;
debug_data->num_bvh_intersections += isect->num_intersections;
debug_data->num_bvh_traversed_nodes += isect.num_traversed_nodes;
debug_data->num_bvh_traversed_instances += isect.num_traversed_instances;
debug_data->num_bvh_intersections += isect.num_intersections;
}
debug_data->num_ray_bounces++;
#endif

4
intern/cycles/kernel/split/kernel_shader_eval.h
View File

@@ -76,14 +76,14 @@ ccl_device void kernel_shader_eval(KernelGlobals *kg)
/* Continue on with shader evaluation. */
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) {
Intersection *isect = &kernel_split_state.isect[ray_index];
Intersection isect = kernel_split_state.isect[ray_index];
ccl_global uint *rng = &kernel_split_state.rng[ray_index];
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
Ray ray = kernel_split_state.ray[ray_index];
shader_setup_from_ray(kg,
&kernel_split_state.sd[ray_index],
isect,
&isect,
&ray);
float rbsdf = path_state_rng_1D_for_decision(kg, rng, state, PRNG_BSDF);
shader_eval_surface(kg, &kernel_split_state.sd[ray_index], rng, state, rbsdf, state->flag, SHADER_CONTEXT_MAIN);

46
intern/cycles/kernel/split/kernel_shadow_blocked.h → intern/cycles/kernel/split/kernel_shadow_blocked_ao.h
View File

@@ -36,42 +36,28 @@ CCL_NAMESPACE_BEGIN
*
* Note on sd_shadow : sd_shadow is neither input nor output to this kernel. sd_shadow is filled and consumed in this kernel itself.
* Note on queues :
* The kernel fetches from QUEUE_SHADOW_RAY_CAST_AO_RAYS and QUEUE_SHADOW_RAY_CAST_DL_RAYS queues. We will empty
* these queues this kernel.
* The kernel fetches from QUEUE_SHADOW_RAY_CAST_AO_RAYS queue. We will empty this queues in this kernel.
* State of queues when this kernel is called :
* state of queues QUEUE_ACTIVE_AND_REGENERATED_RAYS and QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be same
* before and after this kernel call.
* QUEUE_SHADOW_RAY_CAST_AO_RAYS & QUEUE_SHADOW_RAY_CAST_DL_RAYS will be filled with rays marked with flags RAY_SHADOW_RAY_CAST_AO
* and RAY_SHADOW_RAY_CAST_DL respectively, during kernel entry.
* QUEUE_SHADOW_RAY_CAST_AO_RAYS and QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty at kernel exit.
* QUEUE_SHADOW_RAY_CAST_AO_RAYS will be filled with rays marked with flags RAY_SHADOW_RAY_CAST_AO during kernel entry.
* QUEUE_SHADOW_RAY_CAST_AO_RAYS will be empty at kernel exit.
*/
ccl_device void kernel_shadow_blocked(KernelGlobals *kg)
ccl_device void kernel_shadow_blocked_ao(KernelGlobals *kg)
{
int lidx = ccl_local_id(1) * ccl_local_id(0) + ccl_local_id(0);
ccl_local unsigned int ao_queue_length;
ccl_local unsigned int dl_queue_length;
if(lidx == 0) {
ao_queue_length = kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS];
dl_queue_length = kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS];
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
/* flag determining if the current ray is to process shadow ray for AO or DL */
char shadow_blocked_type = -1;
int ray_index = QUEUE_EMPTY_SLOT;
int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
if(thread_index < ao_queue_length + dl_queue_length) {
if(thread_index < ao_queue_length) {
ray_index = get_ray_index(kg, thread_index, QUEUE_SHADOW_RAY_CAST_AO_RAYS,
kernel_split_state.queue_data, kernel_split_params.queue_size, 1);
shadow_blocked_type = RAY_SHADOW_RAY_CAST_AO;
} else {
ray_index = get_ray_index(kg, thread_index - ao_queue_length, QUEUE_SHADOW_RAY_CAST_DL_RAYS,
kernel_split_state.queue_data, kernel_split_params.queue_size, 1);
shadow_blocked_type = RAY_SHADOW_RAY_CAST_DL;
}
if(thread_index < ao_queue_length) {
ray_index = get_ray_index(kg, thread_index, QUEUE_SHADOW_RAY_CAST_AO_RAYS,
kernel_split_state.queue_data, kernel_split_params.queue_size, 1);
}
if(ray_index == QUEUE_EMPTY_SLOT)
@@ -80,25 +66,19 @@ ccl_device void kernel_shadow_blocked(KernelGlobals *kg)
/* Flag determining if we need to update L. */
char update_path_radiance = 0;
if(IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_SHADOW_RAY_CAST_DL) ||
IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_SHADOW_RAY_CAST_AO))
{
if(IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_SHADOW_RAY_CAST_AO)) {
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
ccl_global Ray *light_ray_dl_global = &kernel_split_state.light_ray[ray_index];
ccl_global Ray *light_ray_ao_global = &kernel_split_state.ao_light_ray[ray_index];
ccl_global Ray *light_ray_global =
shadow_blocked_type == RAY_SHADOW_RAY_CAST_AO
? light_ray_ao_global
: light_ray_dl_global;
ccl_global Ray *light_ray_global = &kernel_split_state.ao_light_ray[ray_index];
float3 shadow;
Ray ray = *light_ray_global;
update_path_radiance = !(shadow_blocked(kg,
&kernel_split_state.sd_DL_shadow[thread_index],
&kernel_split_state.sd_DL_shadow[ray_index],
state,
light_ray_global,
&ray,
&shadow));
*light_ray_global = ray;
/* We use light_ray_global's P and t to store shadow and
* update_path_radiance.
*/

91
intern/cycles/kernel/split/kernel_shadow_blocked_dl.h Normal file
View File

@@ -0,0 +1,91 @@
/*
* Copyright 2011-2015 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
/* Note on kernel_shadow_blocked kernel.
* This is the ninth kernel in the ray tracing logic. This is the eighth
* of the path iteration kernels. This kernel takes care of "shadow ray cast"
* logic of the direct lighting and AO part of ray tracing.
*
* The input and output are as follows,
*
* PathState_coop ----------------------------------|--- kernel_shadow_blocked --|
* LightRay_dl_coop --------------------------------| |--- LightRay_dl_coop
* LightRay_ao_coop --------------------------------| |--- LightRay_ao_coop
* ray_state ---------------------------------------| |--- ray_state
* Queue_data(QUEUE_SHADOW_RAY_CAST_AO_RAYS & | |--- Queue_data (QUEUE_SHADOW_RAY_CAST_AO_RAYS & QUEUE_SHADOW_RAY_CAST_AO_RAYS)
QUEUE_SHADOW_RAY_CAST_DL_RAYS) -------| |
* Queue_index(QUEUE_SHADOW_RAY_CAST_AO_RAYS&
QUEUE_SHADOW_RAY_CAST_DL_RAYS) -------| |
* kg (globals) ------------------------------------| |
* queuesize ---------------------------------------| |
*
* Note on sd_shadow : sd_shadow is neither input nor output to this kernel. sd_shadow is filled and consumed in this kernel itself.
* Note on queues :
* The kernel fetches from QUEUE_SHADOW_RAY_CAST_DL_RAYS queue. We will empty this queue in this kernel.
* State of queues when this kernel is called :
* state of queues QUEUE_ACTIVE_AND_REGENERATED_RAYS and QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be same
* before and after this kernel call.
* QUEUE_SHADOW_RAY_CAST_DL_RAYS will be filled with rays marked with flags RAY_SHADOW_RAY_CAST_DL, during kernel entry.
* QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty at kernel exit.
*/
ccl_device void kernel_shadow_blocked_dl(KernelGlobals *kg)
{
int lidx = ccl_local_id(1) * ccl_local_id(0) + ccl_local_id(0);
ccl_local unsigned int dl_queue_length;
if(lidx == 0) {
dl_queue_length = kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS];
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
int ray_index = QUEUE_EMPTY_SLOT;
int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
if(thread_index < dl_queue_length) {
ray_index = get_ray_index(kg, thread_index, QUEUE_SHADOW_RAY_CAST_DL_RAYS,
kernel_split_state.queue_data, kernel_split_params.queue_size, 1);
}
if(ray_index == QUEUE_EMPTY_SLOT)
return;
/* Flag determining if we need to update L. */
char update_path_radiance = 0;
if(IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_SHADOW_RAY_CAST_DL)) {
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
ccl_global Ray *light_ray_global = &kernel_split_state.light_ray[ray_index];
float3 shadow;
Ray ray = *light_ray_global;
update_path_radiance = !(shadow_blocked(kg,
&kernel_split_state.sd_DL_shadow[ray_index],
state,
&ray,
&shadow));
*light_ray_global = ray;
/* We use light_ray_global's P and t to store shadow and
* update_path_radiance.
*/
light_ray_global->P = shadow;
light_ray_global->t = update_path_radiance;
}
}
CCL_NAMESPACE_END

7
intern/cycles/kernel/split/kernel_split_common.h
View File

@@ -52,11 +52,11 @@
#include "kernel_passes.h"
#ifdef __SUBSURFACE__
#include "kernel_subsurface.h"
# include "kernel_subsurface.h"
#endif
#ifdef __VOLUME__
#include "kernel_volume.h"
# include "kernel_volume.h"
#endif
#include "kernel_path_state.h"
@@ -65,9 +65,10 @@
#include "kernel_path_common.h"
#include "kernel_path_surface.h"
#include "kernel_path_volume.h"
#include "kernel_path_subsurface.h"
#ifdef __KERNEL_DEBUG__
#include "kernel_debug.h"
# include "kernel_debug.h"
#endif
#include "kernel_queues.h"

20
intern/cycles/kernel/split/kernel_split_data.h
View File

@@ -31,6 +31,14 @@ ccl_device_inline size_t split_data_buffer_size(KernelGlobals *kg, size_t num_el
size = size SPLIT_DATA_ENTRIES;
#undef SPLIT_DATA_ENTRY
#ifdef __SUBSURFACE__
size += align_up(num_elements * sizeof(SubsurfaceIndirectRays), 16); /* ss_rays */
#endif
#ifdef __VOLUME__
size += align_up(2 * num_elements * sizeof(PathState), 16); /* state_shadow */
#endif
return size;
}
@@ -46,9 +54,19 @@ ccl_device_inline void split_data_init(KernelGlobals *kg,
#define SPLIT_DATA_ENTRY(type, name, num) \
split_data->name = (type*)p; p += align_up(num_elements * num * sizeof(type), 16);
SPLIT_DATA_ENTRIES
SPLIT_DATA_ENTRIES;
#undef SPLIT_DATA_ENTRY
#ifdef __SUBSURFACE__
split_data->ss_rays = (ccl_global SubsurfaceIndirectRays*)p;
p += align_up(num_elements * sizeof(SubsurfaceIndirectRays), 16);
#endif
#ifdef __VOLUME__
split_data->state_shadow = (ccl_global PathState*)p;
p += align_up(2 * num_elements * sizeof(PathState), 16);
#endif
split_data->ray_state = ray_state;
}

13
intern/cycles/kernel/split/kernel_split_data_types.h
View File

@@ -68,18 +68,17 @@ typedef struct SplitParams {
SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \
SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \
SPLIT_DATA_ENTRY(Intersection, isect, 1) \
SPLIT_DATA_ENTRY(ccl_global Intersection, isect, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_alpha, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_bsdf, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ao_light_ray, 1) \
SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \
SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \
SPLIT_DATA_ENTRY(Intersection, isect_shadow, 2) \
SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \
SPLIT_DATA_ENTRY(ccl_global uint, work_array, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 2) \
SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 1) \
SPLIT_DATA_DEBUG_ENTRIES \
/* struct that holds pointers to data in the shared state buffer */
@@ -88,6 +87,14 @@ typedef struct SplitData {
SPLIT_DATA_ENTRIES
#undef SPLIT_DATA_ENTRY
#ifdef __SUBSURFACE__
ccl_global SubsurfaceIndirectRays *ss_rays;
#endif
#ifdef __VOLUME__
ccl_global PathState *state_shadow;
#endif
/* this is actually in a separate buffer from the rest of the split state data (so it can be read back from
* the host easily) but is still used the same as the other data so we have it here in this struct as well
*/

86
intern/cycles/kernel/split/kernel_subsurface_scatter.h Normal file
View File

@@ -0,0 +1,86 @@
CCL_NAMESPACE_BEGIN
ccl_device void kernel_subsurface_scatter(KernelGlobals *kg)
{
#ifdef __SUBSURFACE__
ccl_local unsigned int local_queue_atomics;
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
local_queue_atomics = 0;
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
ray_index = get_ray_index(kg, ray_index,
QUEUE_ACTIVE_AND_REGENERATED_RAYS,
kernel_split_state.queue_data,
kernel_split_params.queue_size,
0);
#ifdef __COMPUTE_DEVICE_GPU__
/* If we are executing on a GPU device, we exit all threads that are not
* required.
*
* If we are executing on a CPU device, then we need to keep all threads
* active since we have barrier() calls later in the kernel. CPU devices,
* expect all threads to execute barrier statement.
*/
if(ray_index == QUEUE_EMPTY_SLOT) {
return;
}
#endif
#ifndef __COMPUTE_DEVICE_GPU__
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
char enqueue_flag = 0;
ccl_global char *ray_state = kernel_split_state.ray_state;
ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
ccl_global RNG *rng = &kernel_split_state.rng[ray_index];
ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
ccl_global SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index];
ShaderData *sd = &kernel_split_state.sd[ray_index];
ShaderData *emission_sd = &kernel_split_state.sd_DL_shadow[ray_index];
if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) {
if(sd->flag & SD_BSSRDF) {
if(kernel_path_subsurface_scatter(kg,
sd,
emission_sd,
L,
state,
rng,
ray,
throughput,
ss_indirect)) {
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
enqueue_flag = 1;
}
}
}
#ifndef __COMPUTE_DEVICE_GPU__
}
#endif
/* Enqueue RAY_UPDATE_BUFFER rays. */
enqueue_ray_index_local(ray_index,
QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,
enqueue_flag,
kernel_split_params.queue_size,
&local_queue_atomics,
kernel_split_state.queue_data,
kernel_split_params.queue_index);
#endif /* __SUBSURFACE__ */
}
CCL_NAMESPACE_END

11
intern/cycles/util/util_debug.cpp
View File

@@ -77,7 +77,8 @@ void DebugFlags::CUDA::reset()
DebugFlags::OpenCL::OpenCL()
: device_type(DebugFlags::OpenCL::DEVICE_ALL),
kernel_type(DebugFlags::OpenCL::KERNEL_DEFAULT),
debug(false)
debug(false),
single_program(false)
{
reset();
}
@@ -117,6 +118,7 @@ void DebugFlags::OpenCL::reset()
}
/* Initialize other flags from environment variables. */
debug = (getenv("CYCLES_OPENCL_DEBUG") != NULL);
single_program = (getenv("CYCLES_OPENCL_SINGLE_PROGRAM") != NULL);
}
DebugFlags::DebugFlags()
@@ -179,9 +181,10 @@ std::ostream& operator <<(std::ostream &os,
break;
}
os << "OpenCL flags:\n"
<< " Device type : " << opencl_device_type << "\n"
<< " Kernel type : " << opencl_kernel_type << "\n"
<< " Debug : " << string_from_bool(debug_flags.opencl.debug)
<< " Device type : " << opencl_device_type << "\n"
<< " Kernel type : " << opencl_kernel_type << "\n"
<< " Debug : " << string_from_bool(debug_flags.opencl.debug) << "\n"
<< " Signle program : " << string_from_bool(debug_flags.opencl.single_program)
<< "\n";
return os;
}

3
intern/cycles/util/util_debug.h
View File

@@ -112,6 +112,9 @@ public:
/* Use debug version of the kernel. */
bool debug;
/* Use single program */
bool single_program;
};
/* Get instance of debug flags registry. */

6
source/blender/editors/armature/armature_select.c
View File

@@ -53,8 +53,6 @@
#include "ED_screen.h"
#include "ED_view3d.h"
#include "GPU_select.h"
#include "armature_intern.h"
/* utility macros for storing a temp int in the bone (selection flag) */
@@ -343,7 +341,7 @@ static EditBone *get_nearest_editbonepoint(
int hits = 0;
/* we _must_ end cache before return, use 'goto cache_end' */
GPU_select_cache_begin();
view3d_opengl_select_cache_begin();
BLI_rcti_init_pt_radius(&rect, mval, 12);
hits12 = view3d_opengl_select(vc, buffer, MAXPICKBUF, &rect, select_mode);
@@ -368,7 +366,7 @@ static EditBone *get_nearest_editbonepoint(
}
cache_end:
GPU_select_cache_end();
view3d_opengl_select_cache_end();
/* See if there are any selected bones in this group */
if (hits > 0) {

10
source/blender/editors/include/ED_view3d.h
View File

@@ -281,7 +281,6 @@ float ED_view3d_radius_to_dist(
const char persp, const bool use_aspect,
const float radius);
void drawcircball(int mode, const float cent[3], float rad, const float tmat[4][4]);
void imm_drawcircball(const float cent[3], float rad, const float tmat[4][4], unsigned pos);
/* backbuffer select and draw support */
@@ -308,18 +307,21 @@ bool ED_view3d_autodist_depth_seg(struct ARegion *ar, const int mval_sta[2], con
#define MAXPICKELEMS 2500
#define MAXPICKBUF (4 * MAXPICKELEMS)
enum {
typedef enum {
/* all elements in the region, ignore depth */
VIEW3D_SELECT_ALL = 0,
/* pick also depth sorts (only for small regions!) */
VIEW3D_SELECT_PICK_ALL = 1,
/* sorts and only returns visible objects (only for small regions!) */
VIEW3D_SELECT_PICK_NEAREST = 2,
};
} eV3DSelectMode;
void view3d_opengl_select_cache_begin(void);
void view3d_opengl_select_cache_end(void);
int view3d_opengl_select(
struct ViewContext *vc, unsigned int *buffer, unsigned int bufsize, const struct rcti *input,
int select_mode);
eV3DSelectMode select_mode);
/* view3d_select.c */
float ED_view3d_select_dist_px(void);

7
source/blender/editors/space_info/textview.c
View File

@@ -207,7 +207,7 @@ static int console_draw_string(ConsoleDrawContext *cdc, const char *str, int str
if (cdc->sel[0] != cdc->sel[1]) {
console_step_sel(cdc, -initial_offset);
// glColor4ub(255, 0, 0, 96); // debug
/* BLF_color3ub(cdc->font_id, 255, 0, 0); // debug */
console_draw_sel(s, cdc->sel, cdc->xy, len, cdc->cwidth, cdc->lheight, bg_sel);
}
@@ -222,7 +222,7 @@ static int console_draw_string(ConsoleDrawContext *cdc, const char *str, int str
if (cdc->sel[0] != cdc->sel[1]) {
console_step_sel(cdc, len);
// glColor4ub(0, 255, 0, 96); // debug
/* BLF_color3ub(cdc->font_id, 0, 255, 0); // debug */
console_draw_sel(s, cdc->sel, cdc->xy, len, cdc->cwidth, cdc->lheight, bg_sel);
}
@@ -251,6 +251,7 @@ static int console_draw_string(ConsoleDrawContext *cdc, const char *str, int str
immUnbindProgram();
}
BLF_color3ubv(cdc->font_id, fg);
BLF_position(cdc->font_id, cdc->xy[0], cdc->lofs + cdc->xy[1], 0);
BLF_draw_mono(cdc->font_id, str, str_len, cdc->cwidth);
@@ -260,7 +261,7 @@ static int console_draw_string(ConsoleDrawContext *cdc, const char *str, int str
isel[0] = str_len - cdc->sel[1];
isel[1] = str_len - cdc->sel[0];
// glColor4ub(255, 255, 0, 96); // debug
/* BLF_color3ub(cdc->font_id, 255, 255, 0); // debug */
console_draw_sel(str, isel, cdc->xy, str_len, cdc->cwidth, cdc->lheight, bg_sel);
console_step_sel(cdc, -(str_len + 1));
}

251
source/blender/editors/space_text/text_draw.c
View File

@@ -49,6 +49,8 @@
#include "BIF_gl.h"
#include "BIF_glutil.h"
#include "GPU_immediate.h"
#include "UI_interface.h"
#include "UI_resources.h"
#include "UI_view2d.h"
@@ -123,38 +125,38 @@ static void txt_format_text(SpaceText *st)
#endif
/* Sets the current drawing color based on the format character specified */
static void format_draw_color(char formatchar)
static void format_draw_color(const TextDrawContext *tdc, char formatchar)
{
switch (formatchar) {
case FMT_TYPE_WHITESPACE:
break;
case FMT_TYPE_SYMBOL:
UI_ThemeColor(TH_SYNTAX_S);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_S);
break;
case FMT_TYPE_COMMENT:
UI_ThemeColor(TH_SYNTAX_C);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_C);
break;
case FMT_TYPE_NUMERAL:
UI_ThemeColor(TH_SYNTAX_N);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_N);
break;
case FMT_TYPE_STRING:
UI_ThemeColor(TH_SYNTAX_L);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_L);
break;
case FMT_TYPE_DIRECTIVE:
UI_ThemeColor(TH_SYNTAX_D);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_D);
break;
case FMT_TYPE_SPECIAL:
UI_ThemeColor(TH_SYNTAX_V);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_V);
break;
case FMT_TYPE_RESERVED:
UI_ThemeColor(TH_SYNTAX_R);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_R);
break;
case FMT_TYPE_KEYWORD:
UI_ThemeColor(TH_SYNTAX_B);
UI_FontThemeColor(tdc->font_id, TH_SYNTAX_B);
break;
case FMT_TYPE_DEFAULT:
default:
UI_ThemeColor(TH_TEXT);
UI_FontThemeColor(tdc->font_id, TH_TEXT);
break;
}
}
@@ -429,7 +431,7 @@ static int text_draw_wrapped(
/* Draw the visible portion of text on the overshot line */
for (a = fstart, ma = mstart; ma < mend; a++, ma += BLI_str_utf8_size_safe(str + ma)) {
if (use_syntax) {
if (fmt_prev != format[a]) format_draw_color(fmt_prev = format[a]);
if (fmt_prev != format[a]) format_draw_color(tdc, fmt_prev = format[a]);
}
x += text_font_draw_character_utf8(tdc, x, y, str + ma);
fpos++;
@@ -452,7 +454,7 @@ static int text_draw_wrapped(
/* Draw the remaining text */
for (a = fstart, ma = mstart; str[ma] && y > clip_min_y; a++, ma += BLI_str_utf8_size_safe(str + ma)) {
if (use_syntax) {
if (fmt_prev != format[a]) format_draw_color(fmt_prev = format[a]);
if (fmt_prev != format[a]) format_draw_color(tdc, fmt_prev = format[a]);
}
x += text_font_draw_character_utf8(tdc, x, y, str + ma);
@@ -505,7 +507,7 @@ static void text_draw(
char fmt_prev = 0xff;
for (a = 0; a < amount; a++) {
if (format[a] != fmt_prev) format_draw_color(fmt_prev = format[a]);
if (format[a] != fmt_prev) format_draw_color(tdc, fmt_prev = format[a]);
x += text_font_draw_character_utf8(tdc, x, y, in + str_shift);
str_shift += BLI_str_utf8_size_safe(in + str_shift);
}
@@ -908,9 +910,13 @@ static void draw_textscroll(const SpaceText *st, rcti *scroll, rcti *back)
uiWidgetColors wcol = btheme->tui.wcol_scroll;
float col[4];
float rad;
UI_ThemeColor(TH_BACK);
glRecti(back->xmin, back->ymin, back->xmax, back->ymax);
/* background so highlights don't go behind the scrollbar */
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
immUniformThemeColor(TH_BACK);
immRecti(pos, back->xmin, back->ymin, back->xmax, back->ymax);
immUnbindProgram();
UI_draw_widget_scroll(&wcol, scroll, &st->txtbar, (st->flags & ST_SCROLL_SELECT) ? UI_SCROLL_PRESSED : 0);
@@ -925,6 +931,7 @@ static void draw_textscroll(const SpaceText *st, rcti *scroll, rcti *back)
/*********************** draw documentation *******************************/
#if 0
static void draw_documentation(const SpaceText *st, ARegion *ar)
{
TextDrawContext tdc = {0};
@@ -961,26 +968,32 @@ static void draw_documentation(const SpaceText *st, ARegion *ar)
boxh = (DOC_HEIGHT + 1) * (st->lheight_dpi + TXT_LINE_SPACING);
/* Draw panel */
UI_ThemeColor(TH_BACK);
glRecti(x, y, x + boxw, y - boxh);
UI_ThemeColor(TH_SHADE1);
glBegin(GL_LINE_LOOP);
glVertex2i(x, y);
glVertex2i(x + boxw, y);
glVertex2i(x + boxw, y - boxh);
glVertex2i(x, y - boxh);
glEnd();
glBegin(GL_LINE_LOOP);
glVertex2i(x + boxw - 10, y - 7);
glVertex2i(x + boxw - 4, y - 7);
glVertex2i(x + boxw - 7, y - 2);
glEnd();
glBegin(GL_LINE_LOOP);
glVertex2i(x + boxw - 10, y - boxh + 7);
glVertex2i(x + boxw - 4, y - boxh + 7);
glVertex2i(x + boxw - 7, y - boxh + 2);
glEnd();
UI_ThemeColor(TH_TEXT);
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
immUniformThemeColor(TH_BACK);
immRecti(pos, x, y, x + boxw, y - boxh);
immUniformThemeColor(TH_SHADE1);
immBegin(GL_LINE_LOOP, 4);
immVertex2i(pos, x, y);
immVertex2i(pos, x + boxw, y);
immVertex2i(pos, x + boxw, y - boxh);
immVertex2i(pos, x, y - boxh);
immEnd();
immBegin(GL_LINE_LOOP, 3);
immVertex2i(pos, x + boxw - 10, y - 7);
immVertex2i(pos, x + boxw - 4, y - 7);
immVertex2i(pos, x + boxw - 7, y - 2);
immEnd();
immBegin(GL_LINE_LOOP, 3);
immVertex2i(pos, x + boxw - 10, y - boxh + 7);
immVertex2i(pos, x + boxw - 4, y - boxh + 7);
immVertex2i(pos, x + boxw - 7, y - boxh + 2);
immEnd();
immUnbindProgram();
UI_FontThemeColor(tdc.font_id, TH_TEXT);
i = 0; br = DOC_WIDTH; lines = 0; // XXX -doc_scroll;
for (p = docs; *p; p++) {
@@ -1007,12 +1020,8 @@ static void draw_documentation(const SpaceText *st, ARegion *ar)
}
if (lines >= DOC_HEIGHT) break;
}
if (0 /* XXX doc_scroll*/ /* > 0 && lines < DOC_HEIGHT */) {
// XXX doc_scroll--;
draw_documentation(st, ar);
}
}
#endif
/*********************** draw suggestion list *******************************/
@@ -1058,10 +1067,15 @@ static void draw_suggestion_list(const SpaceText *st, const TextDrawContext *tdc
/* not needed but stands out nicer */
UI_draw_box_shadow(220, x, y - boxh, x + boxw, y);
UI_ThemeColor(TH_SHADE1);
glRecti(x - 1, y + 1, x + boxw + 1, y - boxh - 1);
UI_ThemeColorShade(TH_BACK, 16);
glRecti(x, y, x + boxw, y - boxh);
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
immUniformThemeColor(TH_SHADE1);
immRecti(pos, x - 1, y + 1, x + boxw + 1, y - boxh - 1);
immUniformThemeColorShade(TH_BACK, 16);
immRecti(pos, x, y, x + boxw, y - boxh);
immUnbindProgram();
/* Set the top 'item' of the visible list */
for (i = 0, item = first; i < *top && item->next; i++, item = item->next) ;
@@ -1076,11 +1090,16 @@ static void draw_suggestion_list(const SpaceText *st, const TextDrawContext *tdc
w = st->cwidth * text_get_char_pos(st, str, len);
if (item == sel) {
UI_ThemeColor(TH_SHADE2);
glRecti(x + margin_x, y - 3, x + margin_x + w, y + lheight - 3);
unsigned int posi = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
immUniformThemeColor(TH_SHADE2);
immRecti(posi, x + margin_x, y - 3, x + margin_x + w, y + lheight - 3);
immUnbindProgram();
}
format_draw_color(item->type);
format_draw_color(tdc, item->type);
text_draw(st, tdc, str, 0, 0, x + margin_x, y - 1, NULL);
if (item == last) break;
@@ -1089,42 +1108,57 @@ static void draw_suggestion_list(const SpaceText *st, const TextDrawContext *tdc
/*********************** draw cursor ************************/
static void draw_cursor(SpaceText *st, ARegion *ar)
static void draw_text_decoration(SpaceText *st, ARegion *ar)
{
Text *text = st->text;
int vcurl, vcurc, vsell, vselc, hidden = 0;
int x, y, w, i;
int offl, offc;
const int lheight = st->lheight_dpi + TXT_LINE_SPACING;
/* Convert to view space character coordinates to determine if cursor is hidden */
wrap_offset(st, ar, text->sell, text->selc, &offl, &offc);
vsell = txt_get_span(text->lines.first, text->sell) - st->top + offl;
vselc = text_get_char_pos(st, text->sell->line, text->selc) - st->left + offc;
if (vselc < 0) {
vselc = 0;
hidden = 1;
}
if (text->curl == text->sell && text->curc == text->selc && !st->line_hlight && hidden) {
/* Nothing to draw here */
return;
}
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
/* Draw the selection */
if (text->curl != text->sell || text->curc != text->selc) {
int offl, offc;
/* Convert all to view space character coordinates */
wrap_offset(st, ar, text->curl, text->curc, &offl, &offc);
vcurl = txt_get_span(text->lines.first, text->curl) - st->top + offl;
vcurc = text_get_char_pos(st, text->curl->line, text->curc) - st->left + offc;
wrap_offset(st, ar, text->sell, text->selc, &offl, &offc);
vsell = txt_get_span(text->lines.first, text->sell) - st->top + offl;
vselc = text_get_char_pos(st, text->sell->line, text->selc) - st->left + offc;
if (vcurc < 0) {
vcurc = 0;
}
if (vselc < 0) {
vselc = 0;
hidden = 1;
}
UI_ThemeColor(TH_SHADE2);
immUniformThemeColor(TH_SHADE2);
x = st->showlinenrs ? TXT_OFFSET + TEXTXLOC : TXT_OFFSET;
y = ar->winy;
if (vcurl == vsell) {
y -= vcurl * lheight;
if (vcurc < vselc)
glRecti(x + vcurc * st->cwidth - 1, y, x + vselc * st->cwidth, y - lheight);
else
glRecti(x + vselc * st->cwidth - 1, y, x + vcurc * st->cwidth, y - lheight);
if (vcurc < vselc) {
immRecti(pos, x + vcurc * st->cwidth - 1, y, x + vselc * st->cwidth, y - lheight);
}
else {
immRecti(pos, x + vselc * st->cwidth - 1, y, x + vcurc * st->cwidth, y - lheight);
}
}
else {
int froml, fromc, tol, toc;
@@ -1140,35 +1174,24 @@ static void draw_cursor(SpaceText *st, ARegion *ar)
y -= froml * lheight;
glRecti(x + fromc * st->cwidth - 1, y, ar->winx, y - lheight);
immRecti(pos, x + fromc * st->cwidth - 1, y, ar->winx, y - lheight);
y -= lheight;
for (i = froml + 1; i < tol; i++) {
glRecti(x - 4, y, ar->winx, y - lheight);
immRecti(pos, x - 4, y, ar->winx, y - lheight);
y -= lheight;
}
glRecti(x - 4, y, x + toc * st->cwidth, y - lheight);
immRecti(pos, x - 4, y, x + toc * st->cwidth, y - lheight);
y -= lheight;
}
}
else {
int offl, offc;
wrap_offset(st, ar, text->sell, text->selc, &offl, &offc);
vsell = txt_get_span(text->lines.first, text->sell) - st->top + offl;
vselc = text_get_char_pos(st, text->sell->line, text->selc) - st->left + offc;
if (vselc < 0) {
vselc = 0;
hidden = 1;
}
}
if (st->line_hlight) {
int x1, x2, y1, y2;
if (st->wordwrap) {
int visible_lines = text_get_visible_lines(st, ar, text->sell->line);
int offl, offc;
wrap_offset_in_line(st, ar, text->sell, text->selc, &offl, &offc);
@@ -1184,36 +1207,38 @@ static void draw_cursor(SpaceText *st, ARegion *ar)
x1 = 0; // st->showlinenrs ? TXT_OFFSET + TEXTXLOC : TXT_OFFSET;
x2 = x1 + ar->winx;
glColor4ub(255, 255, 255, 32);
immUniformColor4ub(255, 255, 255, 32);
glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
glEnable(GL_BLEND);
glRecti(x1 - 4, y1, x2, y2);
immRecti(pos, x1 - 4, y1, x2, y2);
glDisable(GL_BLEND);
}
}
if (!hidden) {
/* Draw the cursor itself (we draw the sel. cursor as this is the leading edge) */
x = st->showlinenrs ? TXT_OFFSET + TEXTXLOC : TXT_OFFSET;
x += vselc * st->cwidth;
y = ar->winy - vsell * lheight;
immUniformThemeColor(TH_HILITE);
if (st->overwrite) {
char ch = text->sell->line[text->selc];
y += TXT_LINE_SPACING;
w = st->cwidth;
if (ch == '\t') w *= st->tabnumber - (vselc + st->left) % st->tabnumber;
UI_ThemeColor(TH_HILITE);
glRecti(x, y - lheight - 1, x + w, y - lheight + 1);
immRecti(pos, x, y - lheight - 1, x + w, y - lheight + 1);
}
else {
UI_ThemeColor(TH_HILITE);
glRecti(x - 1, y, x + 1, y - lheight);
immRecti(pos, x - 1, y, x + 1, y - lheight);
}
}
immUnbindProgram();
}
/******************* draw matching brackets *********************/
@@ -1314,7 +1339,7 @@ static void draw_brackets(const SpaceText *st, const TextDrawContext *tdc, ARegi
if (!endl || endc == -1)
return;
UI_ThemeColor(TH_HILITE);
UI_FontThemeColor(tdc->font_id, TH_HILITE);
x = st->showlinenrs ? TXT_OFFSET + TEXTXLOC : TXT_OFFSET;
y = ar->winy - st->lheight_dpi;
@@ -1419,8 +1444,11 @@ void draw_text_main(SpaceText *st, ARegion *ar)
if (st->showlinenrs) {
x = TXT_OFFSET + TEXTXLOC;
UI_ThemeColor(TH_GRID);
glRecti((TXT_OFFSET - 12), 0, (TXT_OFFSET - 5) + TEXTXLOC, ar->winy - 2);
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
immUniformThemeColor(TH_GRID);
immRecti(pos, (TXT_OFFSET - 12), 0, (TXT_OFFSET - 5) + TEXTXLOC, ar->winy - 2);
immUnbindProgram();
}
else {
st->linenrs_tot = 0; /* not used */
@@ -1429,11 +1457,11 @@ void draw_text_main(SpaceText *st, ARegion *ar)
y = ar->winy - st->lheight_dpi;
winx = ar->winx - TXT_SCROLL_WIDTH;
/* draw cursor */
draw_cursor(st, ar);
/* draw cursor, margin, selection and highlight */
draw_text_decoration(st, ar);
/* draw the text */
UI_ThemeColor(TH_TEXT);
UI_FontThemeColor(tdc.font_id, TH_TEXT);
for (i = 0; y > clip_min_y && i < st->viewlines && tmp; i++, tmp = tmp->next) {
if (st->showsyntax && !tmp->format)
@@ -1441,16 +1469,20 @@ void draw_text_main(SpaceText *st, ARegion *ar)
if (st->showlinenrs && !wrap_skip) {
/* draw line number */
if (tmp == text->curl)
UI_ThemeColor(TH_HILITE);
else
UI_ThemeColor(TH_TEXT);
if (tmp == text->curl) {
UI_FontThemeColor(tdc.font_id, TH_HILITE);
}
else {
UI_FontThemeColor(tdc.font_id, TH_TEXT);
}
BLI_snprintf(linenr, sizeof(linenr), "%*d", st->linenrs_tot, i + linecount + 1);
/* itoa(i + linecount + 1, linenr, 10); */ /* not ansi-c :/ */
text_font_draw(&tdc, TXT_OFFSET - 7, y, linenr);
UI_ThemeColor(TH_TEXT);
if (tmp == text->curl) {
UI_FontThemeColor(tdc.font_id, TH_TEXT);
}
}
if (st->wordwrap) {
@@ -1471,14 +1503,17 @@ void draw_text_main(SpaceText *st, ARegion *ar)
margin_column_x = x + st->cwidth * (st->margin_column - st->left);
if (margin_column_x >= x) {
/* same color as line number background */
UI_ThemeColor(TH_GRID);
setlinestyle(1);
glBegin(GL_LINES);
glVertex2i(margin_column_x, 0);
glVertex2i(margin_column_x, ar->winy - 2);
glEnd();
unsigned int pos = add_attrib(immVertexFormat(), "pos", COMP_I32, 2, CONVERT_INT_TO_FLOAT);
immBindBuiltinProgram(GPU_SHADER_2D_UNIFORM_COLOR);
/* same color as line number background */
immUniformThemeColor(TH_GRID);
immBegin(GL_LINES, 2);
immVertex2i(pos, margin_column_x, 0);
immVertex2i(pos, margin_column_x, ar->winy - 2);
immEnd();
immUnbindProgram();
setlinestyle(0);
}
}
@@ -1486,7 +1521,7 @@ void draw_text_main(SpaceText *st, ARegion *ar)
/* draw other stuff */
draw_brackets(st, &tdc, ar);
draw_textscroll(st, &scroll, &back);
draw_documentation(st, ar);
/* draw_documentation(st, ar); - No longer supported */
draw_suggestion_list(st, &tdc, ar);
text_font_end(&tdc);

711
source/blender/editors/space_view3d/drawobject.c
View File

File diff suppressed because it is too large Load Diff

81
source/blender/editors/space_view3d/view3d_draw_legacy.c
View File

@@ -102,6 +102,7 @@
#include "GPU_compositing.h"
#include "GPU_extensions.h"
#include "GPU_immediate.h"
#include "GPU_select.h"
#include "view3d_intern.h" /* own include */
@@ -908,6 +909,35 @@ static void view3d_draw_xraytransp(Scene *scene, SceneLayer *sl, ARegion *ar, Vi
glDepthMask(GL_TRUE);
}
/* clears zbuffer and draws it over,
* note that in the select version we don't care about transparent flag as with regular drawing */
static void view3d_draw_xray_select(Scene *scene, SceneLayer *sl, ARegion *ar, View3D *v3d, bool *clear)
{
/* Not ideal, but we need to read from the previous depths before clearing
* otherwise we could have a function to load the depths after drawing.
*
* Clearing the depth buffer isn't all that common between drawing objects so accept this for now.
*/
if (U.gpu_select_pick_deph) {
GPU_select_load_id(-1);
}
View3DAfter *v3da;
if (*clear && v3d->zbuf) {
glClear(GL_DEPTH_BUFFER_BIT);
*clear = false;
}
v3d->xray = true;
while ((v3da = BLI_pophead(&v3d->afterdraw_xray))) {
if (GPU_select_load_id(v3da->base->selcol)) {
draw_object_select(scene, sl, ar, v3d, v3da->base, v3da->dflag);
}
MEM_freeN(v3da);
}
v3d->xray = false;
}
/* *********************** */
/*
@@ -1311,6 +1341,57 @@ void ED_view3d_draw_depth(Scene *scene, ARegion *ar, View3D *v3d, bool alphaover
U.obcenter_dia = obcenter_dia;
}
void ED_view3d_draw_select_loop(
ViewContext *vc, Scene *scene, SceneLayer *sl, View3D *v3d, ARegion *ar,
bool use_obedit_skip, bool use_nearest)
{
short code = 1;
const short dflag = DRAW_PICKING | DRAW_CONSTCOLOR;
if (vc->obedit && vc->obedit->type == OB_MBALL) {
draw_object(scene, sl, ar, v3d, BASACT_NEW, dflag);
}
else if ((vc->obedit && vc->obedit->type == OB_ARMATURE)) {
/* if not drawing sketch, draw bones */
if (!BDR_drawSketchNames(vc)) {
draw_object(scene, sl, ar, v3d, BASACT_NEW, dflag);
}
}
else {
Base *base;
for (base = sl->object_bases.first; base; base = base->next) {
if ((base->flag & BASE_VISIBLED) != 0) {
if (((base->flag & BASE_SELECTABLED) == 0) ||
(use_obedit_skip && (scene->obedit->data == base->object->data)))
{
base->selcol = 0;
}
else {
base->selcol = code;
if (use_nearest && (base->object->dtx & OB_DRAWXRAY)) {
ED_view3d_after_add(&v3d->afterdraw_xray, base, dflag);
}
else {
if (GPU_select_load_id(code)) {
draw_object(scene, sl, ar, v3d, base, dflag);
}
}
code++;
}
}
}
if (use_nearest) {
bool xrayclear = true;
if (v3d->afterdraw_xray.first) {
view3d_draw_xray_select(scene, sl, ar, v3d, &xrayclear);
}
}
}
}
typedef struct View3DShadow {
struct View3DShadow *next, *prev;
GPULamp *lamp;

6
source/blender/editors/space_view3d/view3d_intern.h
View File

@@ -145,6 +145,8 @@ void draw_motion_paths_cleanup(View3D *v3d);
/* drawobject.c */
void draw_object(Scene *scene, struct SceneLayer *sl, struct ARegion *ar, View3D *v3d, BaseLegacy *base, const short dflag);
void draw_object_select(Scene *scene, struct SceneLayer *sl, struct ARegion *ar, View3D *v3d, Base *base, const short dflag);
void draw_mesh_object_outline(View3D *v3d, Object *ob, struct DerivedMesh *dm);
bool draw_glsl_material(Scene *scene, struct SceneLayer *sl, struct Object *ob, View3D *v3d, const char dt);
@@ -213,6 +215,10 @@ void view3d_draw_region_info(const struct bContext *C, struct ARegion *ar);
void view3d_main_region_draw_legacy(const struct bContext *C, struct ARegion *ar);
void ED_view3d_draw_depth(Scene *scene, struct ARegion *ar, View3D *v3d, bool alphaoverride);
void ED_view3d_draw_depth_gpencil(Scene *scene, ARegion *ar, View3D *v3d);
void ED_view3d_draw_select_loop(
ViewContext *vc, Scene *scene, struct SceneLayer *sl, View3D *v3d, ARegion *ar,
bool use_obedit_skip, bool use_nearest);
void ED_view3d_after_add(ListBase *lb, BaseLegacy *base, const short dflag);
void circ(float x, float y, float rad);

6
source/blender/editors/space_view3d/view3d_select.c
View File

@@ -98,8 +98,6 @@
#include "GPU_draw.h"
#include "GPU_select.h"
#include "view3d_intern.h" /* own include */
// #include "PIL_time_utildefines.h"
@@ -1216,7 +1214,7 @@ static int mixed_bones_object_selectbuffer(
int hits = 0;
/* we _must_ end cache before return, use 'goto finally' */
GPU_select_cache_begin();
view3d_opengl_select_cache_begin();
BLI_rcti_init_pt_radius(&rect, mval, 14);
hits15 = view3d_opengl_select(vc, buffer, MAXPICKBUF, &rect, select_mode);
@@ -1260,7 +1258,7 @@ static int mixed_bones_object_selectbuffer(
}
finally:
GPU_select_cache_end();
view3d_opengl_select_cache_end();
return hits;
}

85
source/blender/editors/space_view3d/view3d_view.c
View File

@@ -1086,75 +1086,19 @@ void view3d_viewmatrix_set(Scene *scene, const View3D *v3d, RegionView3D *rv3d)
}
}
static void view3d_select_loop(ViewContext *vc, Scene *scene, SceneLayer *sl, View3D *v3d, ARegion *ar, bool use_obedit_skip)
/**
* Optionally cache data for multiple calls to #view3d_opengl_select
*
* just avoid GPU_select headers outside this file
*/
void view3d_opengl_select_cache_begin(void)
{
short code = 1;
char dt;
short dtx;
GPU_select_cache_begin();
}
if (vc->obedit && vc->obedit->type == OB_MBALL) {
draw_object(scene, sl, ar, v3d, BASACT_NEW, DRAW_PICKING | DRAW_CONSTCOLOR);
}
else if ((vc->obedit && vc->obedit->type == OB_ARMATURE)) {
/* if not drawing sketch, draw bones */
if (!BDR_drawSketchNames(vc)) {
draw_object(scene, sl, ar, v3d, BASACT_NEW, DRAW_PICKING | DRAW_CONSTCOLOR);
}
}
else {
Base *base;
v3d->xray = true; /* otherwise it postpones drawing */
for (base = sl->object_bases.first; base; base = base->next) {
if ((base->flag & BASE_VISIBLED) != 0) {
if (((base->flag & BASE_SELECTABLED) == 0) ||
(use_obedit_skip && (scene->obedit->data == base->object->data)))
{
base->selcol = 0;
}
else {
base->selcol = code;
if (GPU_select_load_id(code)) {
draw_object(scene, sl, ar, v3d, base, DRAW_PICKING | DRAW_CONSTCOLOR);
/* we draw duplicators for selection too */
if ((base->object->transflag & OB_DUPLI)) {
ListBase *lb;
DupliObject *dob;
Base tbase;
tbase.flag_legacy = OB_FROMDUPLI;
lb = object_duplilist(G.main->eval_ctx, scene, base->object);
for (dob = lb->first; dob; dob = dob->next) {
float omat[4][4];
tbase.object = dob->ob;
copy_m4_m4(omat, dob->ob->obmat);
copy_m4_m4(dob->ob->obmat, dob->mat);
/* extra service: draw the duplicator in drawtype of parent */
/* MIN2 for the drawtype to allow bounding box objects in groups for lods */
dt = tbase.object->dt; tbase.object->dt = MIN2(tbase.object->dt, base->object->dt);
dtx = tbase.object->dtx; tbase.object->dtx = base->object->dtx;
draw_object(scene, sl, ar, v3d, &tbase, DRAW_PICKING | DRAW_CONSTCOLOR);
tbase.object->dt = dt;
tbase.object->dtx = dtx;
copy_m4_m4(dob->ob->obmat, omat);
}
free_object_duplilist(lb);
}
}
code++;
}
}
}
v3d->xray = false; /* restore */
}
void view3d_opengl_select_cache_end(void)
{
GPU_select_cache_end();
}
/**
@@ -1166,7 +1110,7 @@ static void view3d_select_loop(ViewContext *vc, Scene *scene, SceneLayer *sl, Vi
*/
int view3d_opengl_select(
ViewContext *vc, unsigned int *buffer, unsigned int bufsize, const rcti *input,
int select_mode)
eV3DSelectMode select_mode)
{
Scene *scene = vc->scene;
SceneLayer *sl = vc->sl;
@@ -1180,6 +1124,7 @@ int view3d_opengl_select(
(is_pick_select == false) &&
(select_mode == VIEW3D_SELECT_PICK_NEAREST) &&
GPU_select_query_check_active());
const bool use_nearest = (is_pick_select && select_mode == VIEW3D_SELECT_PICK_NEAREST);
char gpu_select_mode;
@@ -1236,7 +1181,7 @@ int view3d_opengl_select(
GPU_select_begin(buffer, bufsize, &rect, gpu_select_mode, 0);
view3d_select_loop(vc, scene, sl, v3d, ar, use_obedit_skip);
ED_view3d_draw_select_loop(vc, scene, sl, v3d, ar, use_obedit_skip, use_nearest);
hits = GPU_select_end();
@@ -1244,7 +1189,7 @@ int view3d_opengl_select(
if (do_passes) {
GPU_select_begin(buffer, bufsize, &rect, GPU_SELECT_NEAREST_SECOND_PASS, hits);
view3d_select_loop(vc, scene, sl, v3d, ar, use_obedit_skip);
ED_view3d_draw_select_loop(vc, scene, sl, v3d, ar, use_obedit_skip, use_nearest);
GPU_select_end();
}

2
source/blender/gpu/gawain/batch.c
View File

@@ -30,7 +30,7 @@ void Batch_init(Batch* batch, PrimitiveType prim_type, VertexBuffer* verts, Elem
{
#if TRUST_NO_ONE
assert(verts != NULL);
assert(prim_type == PRIM_POINTS || prim_type == PRIM_LINES || prim_type == PRIM_TRIANGLES);
// assert(prim_type == PRIM_POINTS || prim_type == PRIM_LINES || prim_type == PRIM_TRIANGLES);
// we will allow other primitive types in a future update
#endif

44
source/blender/gpu/intern/gpu_select_pick.c
View File

@@ -91,6 +91,15 @@ static void rect_subregion_stride_calc(const rcti *src, const rcti *dst, SubRect
r_sub->skip = src_x - dst_x;
}
/**
* Ignore depth clearing as a change,
* only check if its been changed _and_ filled in (ignore clearing since XRAY does this).
*/
BLI_INLINE bool depth_is_filled(const depth_t *prev, const depth_t *curr)
{
return (*prev != *curr) && (*curr != DEPTH_MAX);
}
/* ----------------------------------------------------------------------------
* DepthBufCache
*
@@ -142,17 +151,28 @@ static bool depth_buf_subrect_depth_any(
return false;
}
static bool depth_buf_rect_not_equal(
const DepthBufCache *rect_depth_a, const DepthBufCache *rect_depth_b,
static bool depth_buf_rect_depth_any_filled(
const DepthBufCache *rect_prev, const DepthBufCache *rect_curr,
unsigned int rect_len)
{
#if 0
return memcmp(rect_depth_a->buf, rect_depth_b->buf, rect_len * sizeof(depth_t)) != 0;
#else
const depth_t *prev = rect_prev->buf;
const depth_t *curr = rect_curr->buf;
for (unsigned int i = 0; i < rect_len; i++, curr++, prev++) {
if (depth_is_filled(prev, curr)) {
return true;
}
}
return false;
#endif
}
/**
* Both buffers are the same size, just check if the sub-rect contains any differences.
*/
static bool depth_buf_subrect_not_equal(
static bool depth_buf_subrect_depth_any_filled(
const DepthBufCache *rect_src, const DepthBufCache *rect_dst,
const SubRectStride *sub_rect)
{
@@ -162,7 +182,7 @@ static bool depth_buf_subrect_not_equal(
for (unsigned int i = 0; i < sub_rect->span_len; i++) {
const depth_t *curr_end = curr + sub_rect->span;
for (; curr < curr_end; prev++, curr++) {
if (*prev != *curr) {
if (depth_is_filled(prev, curr)) {
return true;
}
}
@@ -422,8 +442,11 @@ static void gpu_select_load_id_pass_nearest(const DepthBufCache *rect_prev, cons
if (id != SELECT_ID_NONE) {
unsigned int *id_ptr = ps->nearest.rect_id;
/* Check against DEPTH_MAX because XRAY will clear the buffer,
* so previously set values will become unset.
* In this case just leave those id's left as-is. */
#define EVAL_TEST() \
if (*curr != *prev) { \
if (depth_is_filled(prev, curr)) { \
*id_ptr = id; \
} ((void)0)
@@ -472,7 +495,7 @@ bool gpu_select_pick_load_id(unsigned int id)
}
}
else {
if (depth_buf_rect_not_equal(ps->gl.rect_depth, ps->gl.rect_depth_test, rect_len)) {
if (depth_buf_rect_depth_any_filled(ps->gl.rect_depth, ps->gl.rect_depth_test, rect_len)) {
ps->gl.rect_depth_test->id = ps->gl.prev_id;
gpu_select_load_id_pass_nearest(ps->gl.rect_depth, ps->gl.rect_depth_test);
do_pass = true;
@@ -623,16 +646,17 @@ unsigned int gpu_select_pick_end(void)
hits = -1;
}
else {
/* leave sorting up to the caller */
qsort(depth_data, depth_data_len, sizeof(DepthID), depth_cmp);
for (unsigned int i = 0; i < depth_data_len; i++) {
#ifdef DEBUG_PRINT
printf(" hit: %d: depth %u\n", depth_data[i].id, depth_data[i].depth);
printf(" hit: %u: depth %u\n", depth_data[i].id, depth_data[i].depth);
#endif
/* first 3 are dummy values */
g_pick_state.buffer[hits][0] = 1;
g_pick_state.buffer[hits][1] = 0x0;
g_pick_state.buffer[hits][2] = 0x0;
g_pick_state.buffer[hits][1] = 0x0; /* depth_data[i].depth; */ /* unused */
g_pick_state.buffer[hits][2] = 0x0; /* z-far is currently never used. */
g_pick_state.buffer[hits][3] = depth_data[i].id;
hits++;
}
@@ -709,7 +733,7 @@ void gpu_select_pick_cache_load_id(void)
}
}
else {
if (depth_buf_subrect_not_equal(rect_depth, rect_depth->next, &ps->cache.sub_rect)) {
if (depth_buf_subrect_depth_any_filled(rect_depth, rect_depth->next, &ps->cache.sub_rect)) {
gpu_select_load_id_pass_nearest(rect_depth, rect_depth->next);
}
}