Merge branch 'master' into blender2.8

This commit is contained in:
Sergey Sharybin 2017-06-12 15:09:33 +02:00
commit 0f4f4d8754
114 changed files with 1532 additions and 480 deletions

View File

@ -242,7 +242,8 @@ def register_passes(engine, scene, srl):
if crl.pass_debug_bvh_intersections: engine.register_pass(scene, srl, "Debug BVH Intersections", 1, "X", 'VALUE')
if crl.pass_debug_ray_bounces: engine.register_pass(scene, srl, "Debug Ray Bounces", 1, "X", 'VALUE')
if crl.use_denoising and crl.denoising_store_passes:
cscene = scene.cycles
if crl.use_denoising and crl.denoising_store_passes and not cscene.use_progressive_refine:
engine.register_pass(scene, srl, "Denoising Normal", 3, "XYZ", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Albedo", 3, "RGB", 'COLOR')

View File

@ -689,7 +689,11 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
update=devices_update_callback
)
cls.debug_opencl_kernel_single_program = BoolProperty(name="Single Program", default=True, update=devices_update_callback);
cls.debug_opencl_kernel_single_program = BoolProperty(
name="Single Program",
default=True,
update=devices_update_callback,
)
cls.debug_use_opencl_debug = BoolProperty(name="Debug OpenCL", default=False)
@ -1203,6 +1207,7 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
name="Use Denoising",
description="Denoise the rendered image",
default=False,
update=update_render_passes,
)
cls.denoising_diffuse_direct = BoolProperty(
name="Diffuse Direct",

View File

@ -532,17 +532,17 @@ class CyclesRender_PT_layer_passes(CyclesButtonsPanel, Panel):
col.prop(rl, "use_pass_environment")
if context.scene.cycles.feature_set == 'EXPERIMENTAL':
col.separator()
sub = col.column()
sub.active = crl.use_denoising
sub.prop(crl, "denoising_store_passes", text="Denoising")
col.separator()
sub = col.column()
sub.active = crl.use_denoising
sub.prop(crl, "denoising_store_passes", text="Denoising")
if _cycles.with_cycles_debug:
col = layout.column()
col.prop(crl, "pass_debug_bvh_traversed_nodes")
col.prop(crl, "pass_debug_bvh_traversed_instances")
col.prop(crl, "pass_debug_bvh_intersections")
col.prop(crl, "pass_debug_ray_bounces")
col = layout.column()
col.prop(crl, "pass_debug_bvh_traversed_nodes")
col.prop(crl, "pass_debug_bvh_traversed_instances")
col.prop(crl, "pass_debug_bvh_intersections")
col.prop(crl, "pass_debug_ray_bounces")
class CyclesRender_PT_views(CyclesButtonsPanel, Panel):
@ -1688,7 +1688,7 @@ def draw_device(self, context):
layout.prop(cscene, "feature_set")
split = layout.split(percentage=1/3)
split = layout.split(percentage=1 / 3)
split.label("Device:")
row = split.row()
row.active = show_device_active(context)

View File

@ -403,14 +403,7 @@ void BlenderSession::render()
BL::RenderLayer b_rlay = *b_single_rlay;
/* add passes */
array<Pass> passes;
if(session_params.device.advanced_shading) {
passes = sync->sync_render_passes(b_rlay, *b_layer_iter);
}
else {
Pass::add(PASS_COMBINED, passes);
}
array<Pass> passes = sync->sync_render_passes(b_rlay, *b_layer_iter, session_params);
buffer_params.passes = passes;
PointerRNA crl = RNA_pointer_get(&b_layer_iter->ptr, "cycles");

View File

@ -537,11 +537,16 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
}
array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
BL::SceneRenderLayer& b_srlay)
BL::SceneRenderLayer& b_srlay,
const SessionParams &session_params)
{
array<Pass> passes;
Pass::add(PASS_COMBINED, passes);
if(!session_params.device.advanced_shading) {
return passes;
}
/* loop over passes */
BL::RenderLayer::passes_iterator b_pass_iter;
@ -556,7 +561,9 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
}
PointerRNA crp = RNA_pointer_get(&b_srlay.ptr, "cycles");
if(get_boolean(crp, "denoising_store_passes")) {
if(get_boolean(crp, "denoising_store_passes") &&
get_boolean(crp, "use_denoising") &&
!session_params.progressive_refine) {
b_engine.add_pass("Denoising Normal", 3, "XYZ", b_srlay.name().c_str());
b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_srlay.name().c_str());
b_engine.add_pass("Denoising Albedo", 3, "RGB", b_srlay.name().c_str());

View File

@ -69,7 +69,8 @@ public:
const char *layer = 0);
void sync_render_layers(BL::SpaceView3D& b_v3d, const char *layer);
array<Pass> sync_render_passes(BL::RenderLayer& b_rlay,
BL::SceneRenderLayer& b_srlay);
BL::SceneRenderLayer& b_srlay,
const SessionParams &session_params);
void sync_integrator();
void sync_camera(BL::RenderSettings& b_render,
BL::Object& b_override,

View File

@ -69,6 +69,8 @@ std::ostream& operator <<(std::ostream &os,
<< string_from_bool(requested_features.use_transparent) << std::endl;
os << "Use Principled BSDF: "
<< string_from_bool(requested_features.use_principled) << std::endl;
os << "Use Denoising: "
<< string_from_bool(requested_features.use_denoising) << std::endl;
return os;
}

View File

@ -127,6 +127,9 @@ public:
/* Per-uber shader usage flags. */
bool use_principled;
/* Denoising features. */
bool use_denoising;
DeviceRequestedFeatures()
{
/* TODO(sergey): Find more meaningful defaults. */
@ -145,6 +148,7 @@ public:
use_transparent = false;
use_shadow_tricks = false;
use_principled = false;
use_denoising = false;
}
bool modified(const DeviceRequestedFeatures& requested_features)
@ -163,7 +167,8 @@ public:
use_patch_evaluation == requested_features.use_patch_evaluation &&
use_transparent == requested_features.use_transparent &&
use_shadow_tricks == requested_features.use_shadow_tricks &&
use_principled == requested_features.use_principled);
use_principled == requested_features.use_principled &&
use_denoising == requested_features.use_denoising);
}
/* Convert the requested features structure to a build options,
@ -213,6 +218,9 @@ public:
if(!use_principled) {
build_options += " -D__NO_PRINCIPLED__";
}
if(!use_denoising) {
build_options += " -D__NO_DENOISING__";
}
return build_options;
}
};

View File

@ -47,6 +47,7 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device) : device(device)
kernel_direct_lighting = NULL;
kernel_shadow_blocked_ao = NULL;
kernel_shadow_blocked_dl = NULL;
kernel_enqueue_inactive = NULL;
kernel_next_iteration_setup = NULL;
kernel_indirect_subsurface = NULL;
kernel_buffer_update = NULL;
@ -74,6 +75,7 @@ DeviceSplitKernel::~DeviceSplitKernel()
delete kernel_direct_lighting;
delete kernel_shadow_blocked_ao;
delete kernel_shadow_blocked_dl;
delete kernel_enqueue_inactive;
delete kernel_next_iteration_setup;
delete kernel_indirect_subsurface;
delete kernel_buffer_update;
@ -101,6 +103,7 @@ bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures& requested_fe
LOAD_KERNEL(direct_lighting);
LOAD_KERNEL(shadow_blocked_ao);
LOAD_KERNEL(shadow_blocked_dl);
LOAD_KERNEL(enqueue_inactive);
LOAD_KERNEL(next_iteration_setup);
LOAD_KERNEL(indirect_subsurface);
LOAD_KERNEL(buffer_update);
@ -256,6 +259,7 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shadow_blocked_ao, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shadow_blocked_dl, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(enqueue_inactive, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(indirect_subsurface, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size);

View File

@ -69,6 +69,7 @@ private:
SplitKernelFunction *kernel_direct_lighting;
SplitKernelFunction *kernel_shadow_blocked_ao;
SplitKernelFunction *kernel_shadow_blocked_dl;
SplitKernelFunction *kernel_enqueue_inactive;
SplitKernelFunction *kernel_next_iteration_setup;
SplitKernelFunction *kernel_indirect_subsurface;
SplitKernelFunction *kernel_buffer_update;

View File

@ -130,6 +130,11 @@ public:
cl_int* error = NULL);
static cl_device_type get_device_type(cl_device_id device_id);
static bool get_driver_version(cl_device_id device_id,
int *major,
int *minor,
cl_int* error = NULL);
static int mem_address_alignment(cl_device_id device_id);
/* Get somewhat more readable device name.

View File

@ -176,17 +176,62 @@ protected:
friend class OpenCLSplitKernelFunction;
};
struct CachedSplitMemory {
int id;
device_memory *split_data;
device_memory *ray_state;
device_ptr *rng_state;
device_memory *queue_index;
device_memory *use_queues_flag;
device_memory *work_pools;
device_ptr *buffer;
};
class OpenCLSplitKernelFunction : public SplitKernelFunction {
public:
OpenCLDeviceSplitKernel* device;
OpenCLDeviceBase::OpenCLProgram program;
CachedSplitMemory& cached_memory;
int cached_id;
OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device) : device(device) {}
~OpenCLSplitKernelFunction() { program.release(); }
OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device, CachedSplitMemory& cached_memory) :
device(device), cached_memory(cached_memory), cached_id(cached_memory.id-1)
{
}
~OpenCLSplitKernelFunction()
{
program.release();
}
virtual bool enqueue(const KernelDimensions& dim, device_memory& kg, device_memory& data)
{
device->kernel_set_args(program(), 0, kg, data);
if(cached_id != cached_memory.id) {
cl_uint start_arg_index =
device->kernel_set_args(program(),
0,
kg,
data,
*cached_memory.split_data,
*cached_memory.ray_state,
*cached_memory.rng_state);
/* TODO(sergey): Avoid map lookup here. */
#define KERNEL_TEX(type, ttype, name) \
device->set_kernel_arg_mem(program(), &start_arg_index, #name);
#include "kernel/kernel_textures.h"
#undef KERNEL_TEX
start_arg_index +=
device->kernel_set_args(program(),
start_arg_index,
*cached_memory.queue_index,
*cached_memory.use_queues_flag,
*cached_memory.work_pools,
*cached_memory.buffer);
cached_id = cached_memory.id;
}
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
program(),
@ -213,6 +258,7 @@ public:
class OpenCLSplitKernel : public DeviceSplitKernel {
OpenCLDeviceSplitKernel *device;
CachedSplitMemory cached_memory;
public:
explicit OpenCLSplitKernel(OpenCLDeviceSplitKernel *device) : DeviceSplitKernel(device), device(device) {
}
@ -220,7 +266,7 @@ public:
virtual SplitKernelFunction* get_split_kernel_function(string kernel_name,
const DeviceRequestedFeatures& requested_features)
{
OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device);
OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device, cached_memory);
bool single_program = OpenCLInfo::use_single_program();
kernel->program =
@ -349,6 +395,15 @@ public:
return false;
}
cached_memory.split_data = &split_data;
cached_memory.ray_state = &ray_state;
cached_memory.rng_state = &rtile.rng_state;
cached_memory.queue_index = &queue_index;
cached_memory.use_queues_flag = &use_queues_flag;
cached_memory.work_pools = &work_pool_wgs;
cached_memory.buffer = &rtile.buffer;
cached_memory.id++;
return true;
}

View File

@ -608,6 +608,14 @@ bool OpenCLInfo::device_supported(const string& platform_name,
if(!get_device_name(device_id, &device_name)) {
return false;
}
int driver_major = 0;
int driver_minor = 0;
if(!get_driver_version(device_id, &driver_major, &driver_minor)) {
return false;
}
VLOG(3) << "OpenCL driver version " << driver_major << "." << driver_minor;
/* It is possible tyo have Iris GPU on AMD/Apple OpenCL framework
* (aka, it will not be on Intel framework). This isn't supported
* and needs an explicit blacklist.
@ -618,6 +626,21 @@ bool OpenCLInfo::device_supported(const string& platform_name,
if(platform_name == "AMD Accelerated Parallel Processing" &&
device_type == CL_DEVICE_TYPE_GPU)
{
if(driver_major < 2236) {
VLOG(1) << "AMD driver version " << driver_major << "." << driver_minor << " not supported.";
return false;
}
const char *blacklist[] = {
/* GCN 1 */
"Tahiti", "Pitcairn", "Capeverde", "Oland",
NULL
};
for (int i = 0; blacklist[i] != NULL; i++) {
if(device_name == blacklist[i]) {
VLOG(1) << "AMD device " << device_name << " not supported";
return false;
}
}
return true;
}
if(platform_name == "Apple" && device_type == CL_DEVICE_TYPE_GPU) {
@ -1073,6 +1096,34 @@ string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
return get_device_name(device_id);
}
bool OpenCLInfo::get_driver_version(cl_device_id device_id,
int *major,
int *minor,
cl_int* error)
{
char buffer[1024];
cl_int err;
if((err = clGetDeviceInfo(device_id,
CL_DRIVER_VERSION,
sizeof(buffer),
&buffer,
NULL)) != CL_SUCCESS)
{
if(error != NULL) {
*error = err;
}
return false;
}
if(error != NULL) {
*error = CL_SUCCESS;
}
if(sscanf(buffer, "%d.%d", major, minor) < 2) {
VLOG(1) << string_printf("OpenCL: failed to parse driver version string (%s).", buffer);
return false;
}
return true;
}
int OpenCLInfo::mem_address_alignment(cl_device_id device_id)
{
int base_align_bits;

View File

@ -45,6 +45,7 @@ set(SRC
kernels/opencl/kernel_direct_lighting.cl
kernels/opencl/kernel_shadow_blocked_ao.cl
kernels/opencl/kernel_shadow_blocked_dl.cl
kernels/opencl/kernel_enqueue_inactive.cl
kernels/opencl/kernel_next_iteration_setup.cl
kernels/opencl/kernel_indirect_subsurface.cl
kernels/opencl/kernel_buffer_update.cl
@ -121,6 +122,10 @@ set(SRC_KERNELS_CUDA_HEADERS
kernels/cuda/kernel_config.h
)
set(SRC_KERNELS_OPENCL_HEADERS
kernels/opencl/kernel_split_function.h
)
set(SRC_CLOSURE_HEADERS
closure/alloc.h
closure/bsdf.h
@ -278,6 +283,7 @@ set(SRC_SPLIT_HEADERS
split/kernel_data_init.h
split/kernel_direct_lighting.h
split/kernel_do_volume.h
split/kernel_enqueue_inactive.h
split/kernel_holdout_emission_blurring_pathtermination_ao.h
split/kernel_indirect_background.h
split/kernel_indirect_subsurface.h
@ -450,6 +456,7 @@ add_library(cycles_kernel
${SRC_HEADERS}
${SRC_KERNELS_CPU_HEADERS}
${SRC_KERNELS_CUDA_HEADERS}
${SRC_KERNELS_OPENCL_HEADERS}
${SRC_BVH_HEADERS}
${SRC_CLOSURE_HEADERS}
${SRC_FILTER_HEADERS}
@ -490,9 +497,11 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_subsurface_sc
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_direct_lighting.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_ao.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocked_dl.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_enqueue_inactive.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_next_iteration_setup.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_indirect_subsurface.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_buffer_update.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_split_function.h" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/filter.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel_split.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda)

View File

@ -101,7 +101,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict d
for(int x = rect.x; x < rect.z; x++) {
const int low = max(rect.x, x-f);
const int high = min(rect.z, x+f+1);
out_image[y*w+x] = expf(-max(out_image[y*w+x] * (1.0f/(high - low)), 0.0f));
out_image[y*w+x] = fast_expf(-max(out_image[y*w+x] * (1.0f/(high - low)), 0.0f));
}
}
}

View File

@ -66,7 +66,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(int x, int y,
sum += difference_image[y*w+x1];
}
sum *= 1.0f/(high-low);
out_image[y*w+x] = expf(-max(sum, 0.0f));
out_image[y*w+x] = fast_expf(-max(sum, 0.0f));
}
ccl_device_inline void kernel_filter_nlm_update_output(int x, int y,

View File

@ -29,20 +29,24 @@ ccl_device_inline void kernel_filter_construct_gramian(int x, int y,
ccl_global float3 *XtWY,
int localIdx)
{
if(weight < 1e-3f) {
return;
}
int p_offset = y *w + x;
int q_offset = (y+dy)*w + (x+dx);
#ifdef __KERNEL_CPU__
const int stride = 1;
(void)storage_stride;
(void)localIdx;
float design_row[DENOISE_FEATURES+1];
#elif defined(__KERNEL_CUDA__)
#ifdef __KERNEL_GPU__
const int stride = storage_stride;
#else
const int stride = 1;
(void) storage_stride;
#endif
#ifdef __KERNEL_CUDA__
ccl_local float shared_design_row[(DENOISE_FEATURES+1)*CCL_MAX_LOCAL_SIZE];
ccl_local_param float *design_row = shared_design_row + localIdx*(DENOISE_FEATURES+1);
#else
const int stride = storage_stride;
float design_row[DENOISE_FEATURES+1];
#endif
@ -70,13 +74,19 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h,
int4 buffer_params,
int sample)
{
#ifdef __KERNEL_CPU__
const int stride = 1;
(void)storage_stride;
#else
#ifdef __KERNEL_GPU__
const int stride = storage_stride;
#else
const int stride = 1;
(void) storage_stride;
#endif
if(XtWX[0] < 1e-3f) {
/* There is not enough information to determine a denoised result.
* As a fallback, keep the original value of the pixel. */
return;
}
/* The weighted average of pixel colors (essentially, the NLM-filtered image).
* In case the solution of the linear model fails due to numerical issues,
* fall back to this value. */
@ -89,6 +99,9 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h,
final_color = mean_color;
}
/* Clamp pixel value to positive values. */
final_color = max(final_color, make_float3(0.0f, 0.0f, 0.0f));
ccl_global float *combined_buffer = buffer + (y*buffer_params.y + x + buffer_params.x)*buffer_params.z;
final_color *= sample;
if(buffer_params.w) {
@ -101,6 +114,4 @@ ccl_device_inline void kernel_filter_finalize(int x, int y, int w, int h,
combined_buffer[2] = final_color.z;
}
#undef STORAGE_TYPE
CCL_NAMESPACE_END

View File

@ -621,25 +621,43 @@ ccl_device_inline void path_radiance_accum_sample(PathRadiance *L, PathRadiance
{
float fac = 1.0f/num_samples;
#ifdef __SPLIT_KERNEL__
# define safe_float3_add(f, v) \
do { \
ccl_global float *p = (ccl_global float*)(&(f)); \
atomic_add_and_fetch_float(p+0, (v).x); \
atomic_add_and_fetch_float(p+1, (v).y); \
atomic_add_and_fetch_float(p+2, (v).z); \
} while(0)
#else
# define safe_float3_add(f, v) (f) += (v)
#endif /* __SPLIT_KERNEL__ */
#ifdef __PASSES__
L->direct_diffuse += L_sample->direct_diffuse*fac;
L->direct_glossy += L_sample->direct_glossy*fac;
L->direct_transmission += L_sample->direct_transmission*fac;
L->direct_subsurface += L_sample->direct_subsurface*fac;
L->direct_scatter += L_sample->direct_scatter*fac;
safe_float3_add(L->direct_diffuse, L_sample->direct_diffuse*fac);
safe_float3_add(L->direct_glossy, L_sample->direct_glossy*fac);
safe_float3_add(L->direct_transmission, L_sample->direct_transmission*fac);
safe_float3_add(L->direct_subsurface, L_sample->direct_subsurface*fac);
safe_float3_add(L->direct_scatter, L_sample->direct_scatter*fac);
L->indirect_diffuse += L_sample->indirect_diffuse*fac;
L->indirect_glossy += L_sample->indirect_glossy*fac;
L->indirect_transmission += L_sample->indirect_transmission*fac;
L->indirect_subsurface += L_sample->indirect_subsurface*fac;
L->indirect_scatter += L_sample->indirect_scatter*fac;
safe_float3_add(L->indirect_diffuse, L_sample->indirect_diffuse*fac);
safe_float3_add(L->indirect_glossy, L_sample->indirect_glossy*fac);
safe_float3_add(L->indirect_transmission, L_sample->indirect_transmission*fac);
safe_float3_add(L->indirect_subsurface, L_sample->indirect_subsurface*fac);
safe_float3_add(L->indirect_scatter, L_sample->indirect_scatter*fac);
L->background += L_sample->background*fac;
L->ao += L_sample->ao*fac;
L->shadow += L_sample->shadow*fac;
safe_float3_add(L->background, L_sample->background*fac);
safe_float3_add(L->ao, L_sample->ao*fac);
safe_float3_add(L->shadow, L_sample->shadow*fac);
# ifdef __SPLIT_KERNEL__
atomic_add_and_fetch_float(&L->mist, L_sample->mist*fac);
# else
L->mist += L_sample->mist*fac;
#endif
L->emission += L_sample->emission * fac;
# endif /* __SPLIT_KERNEL__ */
#endif /* __PASSES__ */
safe_float3_add(L->emission, L_sample->emission*fac);
#undef safe_float3_add
}
#ifdef __SHADOW_TRICKS__

View File

@ -139,9 +139,11 @@ ccl_device_inline void path_state_next(KernelGlobals *kg, ccl_addr_space PathSta
/* random number generator next bounce */
state->rng_offset += PRNG_BOUNCE_NUM;
#ifdef __DENOISING_FEATURES__
if((state->denoising_feature_weight == 0.0f) && !(state->flag & PATH_RAY_SHADOW_CATCHER)) {
state->flag &= ~PATH_RAY_STORE_SHADOW_INFO;
}
#endif
}
ccl_device_inline uint path_state_ray_visibility(KernelGlobals *kg, PathState *state)

View File

@ -128,6 +128,21 @@ ccl_device unsigned int get_global_queue_index(
return my_gqidx;
}
ccl_device int dequeue_ray_index(
int queue_number,
ccl_global int *queues,
int queue_size,
ccl_global int *queue_index)
{
int index = atomic_fetch_and_dec_uint32((ccl_global uint*)&queue_index[queue_number])-1;
if(index < 0) {
return QUEUE_EMPTY_SLOT;
}
return queues[index + queue_number * queue_size];
}
CCL_NAMESPACE_END
#endif // __KERNEL_QUEUE_H__

View File

@ -236,6 +236,9 @@ CCL_NAMESPACE_BEGIN
#ifdef __NO_PRINCIPLED__
# undef __PRINCIPLED__
#endif
#ifdef __NO_DENOISING__
# undef __DENOISING_FEATURES__
#endif
/* Random Numbers */
@ -1387,6 +1390,8 @@ enum QueueNumber {
#ifdef __BRANCHED_PATH__
/* All rays moving to next iteration of the indirect loop for light */
QUEUE_LIGHT_INDIRECT_ITER,
/* Queue of all inactive rays. These are candidates for sharing work of indirect loops */
QUEUE_INACTIVE_RAYS,
# ifdef __VOLUME__
/* All rays moving to next iteration of the indirect loop for volumes */
QUEUE_VOLUME_INDIRECT_ITER,
@ -1429,6 +1434,9 @@ enum RayState {
RAY_BRANCHED_VOLUME_INDIRECT = (1 << 5),
RAY_BRANCHED_SUBSURFACE_INDIRECT = (1 << 6),
RAY_BRANCHED_INDIRECT = (RAY_BRANCHED_LIGHT_INDIRECT | RAY_BRANCHED_VOLUME_INDIRECT | RAY_BRANCHED_SUBSURFACE_INDIRECT),
/* Ray is evaluating an iteration of an indirect loop for another thread */
RAY_BRANCHED_INDIRECT_SHARED = (1 << 7),
};
#define ASSIGN_RAY_STATE(ray_state, ray_index, state) (ray_state[ray_index] = ((ray_state[ray_index] & RAY_FLAG_MASK) | state))

View File

@ -85,6 +85,7 @@ DECLARE_SPLIT_KERNEL_FUNCTION(subsurface_scatter)
DECLARE_SPLIT_KERNEL_FUNCTION(direct_lighting)
DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao)
DECLARE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl)
DECLARE_SPLIT_KERNEL_FUNCTION(enqueue_inactive)
DECLARE_SPLIT_KERNEL_FUNCTION(next_iteration_setup)
DECLARE_SPLIT_KERNEL_FUNCTION(indirect_subsurface)
DECLARE_SPLIT_KERNEL_FUNCTION(buffer_update)

View File

@ -53,6 +53,7 @@
# include "kernel/split/kernel_direct_lighting.h"
# include "kernel/split/kernel_shadow_blocked_ao.h"
# include "kernel/split/kernel_shadow_blocked_dl.h"
# include "kernel/split/kernel_enqueue_inactive.h"
# include "kernel/split/kernel_next_iteration_setup.h"
# include "kernel/split/kernel_indirect_subsurface.h"
# include "kernel/split/kernel_buffer_update.h"
@ -230,6 +231,7 @@ DEFINE_SPLIT_KERNEL_FUNCTION(subsurface_scatter)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(direct_lighting, uint)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint)
DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint)

View File

@ -39,6 +39,7 @@
#include "kernel/split/kernel_direct_lighting.h"
#include "kernel/split/kernel_shadow_blocked_ao.h"
#include "kernel/split/kernel_shadow_blocked_dl.h"
#include "kernel/split/kernel_enqueue_inactive.h"
#include "kernel/split/kernel_next_iteration_setup.h"
#include "kernel/split/kernel_indirect_subsurface.h"
#include "kernel/split/kernel_buffer_update.h"
@ -118,6 +119,7 @@ DEFINE_SPLIT_KERNEL_FUNCTION(subsurface_scatter)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(direct_lighting, uint)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_ao)
DEFINE_SPLIT_KERNEL_FUNCTION(shadow_blocked_dl)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint)
DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface)
DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint)

View File

@ -18,10 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_buffer_update.h"
__kernel void kernel_ocl_path_trace_buffer_update(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local unsigned int local_queue_atomics;
kernel_buffer_update((KernelGlobals*)kg, &local_queue_atomics);
}
#define KERNEL_NAME buffer_update
#define LOCALS_TYPE unsigned int
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,10 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_direct_lighting.h"
__kernel void kernel_ocl_path_trace_direct_lighting(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local unsigned int local_queue_atomics;
kernel_direct_lighting((KernelGlobals*)kg, &local_queue_atomics);
}
#define KERNEL_NAME direct_lighting
#define LOCALS_TYPE unsigned int
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_do_volume.h"
__kernel void kernel_ocl_path_trace_do_volume(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_do_volume((KernelGlobals*)kg);
}
#define KERNEL_NAME do_volume
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

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/kernel_compat_opencl.h"
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_enqueue_inactive.h"
#define KERNEL_NAME enqueue_inactive
#define LOCALS_TYPE unsigned int
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,12 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h"
__kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local BackgroundAOLocals locals;
kernel_holdout_emission_blurring_pathtermination_ao(
(KernelGlobals*)kg,
&locals);
}
#define KERNEL_NAME holdout_emission_blurring_pathtermination_ao
#define LOCALS_TYPE BackgroundAOLocals
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_indirect_background.h"
__kernel void kernel_ocl_path_trace_indirect_background(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_indirect_background((KernelGlobals*)kg);
}
#define KERNEL_NAME indirect_background
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_indirect_subsurface.h"
__kernel void kernel_ocl_path_trace_indirect_subsurface(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_indirect_subsurface((KernelGlobals*)kg);
}
#define KERNEL_NAME indirect_subsurface
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_lamp_emission.h"
__kernel void kernel_ocl_path_trace_lamp_emission(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_lamp_emission((KernelGlobals*)kg);
}
#define KERNEL_NAME lamp_emission
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,10 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_next_iteration_setup.h"
__kernel void kernel_ocl_path_trace_next_iteration_setup(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local unsigned int local_queue_atomics;
kernel_next_iteration_setup((KernelGlobals*)kg, &local_queue_atomics);
}
#define KERNEL_NAME next_iteration_setup
#define LOCALS_TYPE unsigned int
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_path_init.h"
__kernel void kernel_ocl_path_trace_path_init(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_path_init((KernelGlobals*)kg);
}
#define KERNEL_NAME path_init
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,10 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_queue_enqueue.h"
__kernel void kernel_ocl_path_trace_queue_enqueue(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local QueueEnqueueLocals locals;
kernel_queue_enqueue((KernelGlobals*)kg, &locals);
}
#define KERNEL_NAME queue_enqueue
#define LOCALS_TYPE QueueEnqueueLocals
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_scene_intersect.h"
__kernel void kernel_ocl_path_trace_scene_intersect(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_scene_intersect((KernelGlobals*)kg);
}
#define KERNEL_NAME scene_intersect
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_shader_eval.h"
__kernel void kernel_ocl_path_trace_shader_eval(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_shader_eval((KernelGlobals*)kg);
}
#define KERNEL_NAME shader_eval
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,10 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_shader_setup.h"
__kernel void kernel_ocl_path_trace_shader_setup(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local unsigned int local_queue_atomics;
kernel_shader_setup((KernelGlobals*)kg, &local_queue_atomics);
}
#define KERNEL_NAME shader_setup
#define LOCALS_TYPE unsigned int
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -19,10 +19,9 @@
#include "kernel/split/kernel_shader_sort.h"
__attribute__((reqd_work_group_size(64, 1, 1)))
__kernel void kernel_ocl_path_trace_shader_sort(
ccl_global char *kg,
ccl_constant KernelData *data)
{
ccl_local ShaderSortLocals locals;
kernel_shader_sort((KernelGlobals*)kg, &locals);
}
#define KERNEL_NAME shader_sort
#define LOCALS_TYPE ShaderSortLocals
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME
#undef LOCALS_TYPE

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_shadow_blocked_ao.h"
__kernel void kernel_ocl_path_trace_shadow_blocked_ao(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_shadow_blocked_ao((KernelGlobals*)kg);
}
#define KERNEL_NAME shadow_blocked_ao
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_shadow_blocked_dl.h"
__kernel void kernel_ocl_path_trace_shadow_blocked_dl(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_shadow_blocked_dl((KernelGlobals*)kg);
}
#define KERNEL_NAME shadow_blocked_dl
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -31,6 +31,7 @@
#include "kernel/kernels/opencl/kernel_direct_lighting.cl"
#include "kernel/kernels/opencl/kernel_shadow_blocked_ao.cl"
#include "kernel/kernels/opencl/kernel_shadow_blocked_dl.cl"
#include "kernel/kernels/opencl/kernel_enqueue_inactive.cl"
#include "kernel/kernels/opencl/kernel_next_iteration_setup.cl"
#include "kernel/kernels/opencl/kernel_indirect_subsurface.cl"
#include "kernel/kernels/opencl/kernel_buffer_update.cl"

View File

@ -0,0 +1,72 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#define KERNEL_NAME_JOIN(a, b) a ## _ ## b
#define KERNEL_NAME_EVAL(a, b) KERNEL_NAME_JOIN(a, b)
__kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)(
ccl_global char *kg_global,
ccl_constant KernelData *data,
ccl_global void *split_data_buffer,
ccl_global char *ray_state,
ccl_global uint *rng_state,
#define KERNEL_TEX(type, ttype, name) \
ccl_global type *name,
#include "kernel/kernel_textures.h"
ccl_global int *queue_index,
ccl_global char *use_queues_flag,
ccl_global unsigned int *work_pools,
ccl_global float *buffer
)
{
#ifdef LOCALS_TYPE
ccl_local LOCALS_TYPE locals;
#endif
KernelGlobals *kg = (KernelGlobals*)kg_global;
if(ccl_local_id(0) + ccl_local_id(1) == 0) {
kg->data = data;
kernel_split_params.rng_state = rng_state;
kernel_split_params.queue_index = queue_index;
kernel_split_params.use_queues_flag = use_queues_flag;
kernel_split_params.work_pools = work_pools;
kernel_split_params.buffer = buffer;
split_data_init(kg, &kernel_split_state, ccl_global_size(0)*ccl_global_size(1), split_data_buffer, ray_state);
#define KERNEL_TEX(type, ttype, name) \
kg->name = name;
#include "kernel/kernel_textures.h"
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
KERNEL_NAME_EVAL(kernel, KERNEL_NAME)(
kg
#ifdef LOCALS_TYPE
, &locals
#endif
);
}
#undef KERNEL_NAME_JOIN
#undef KERNEL_NAME_EVAL

View File

@ -18,9 +18,7 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_subsurface_scatter.h"
__kernel void kernel_ocl_path_trace_subsurface_scatter(
ccl_global char *kg,
ccl_constant KernelData *data)
{
kernel_subsurface_scatter((KernelGlobals*)kg);
}
#define KERNEL_NAME subsurface_scatter
#include "kernel/kernels/opencl/kernel_split_function.h"
#undef KERNEL_NAME

View File

@ -63,12 +63,49 @@ ccl_device_inline void kernel_split_branched_path_indirect_loop_end(KernelGlobal
REMOVE_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_INDIRECT);
}
ccl_device_inline bool kernel_split_branched_indirect_start_shared(KernelGlobals *kg, int ray_index)
{
ccl_global char *ray_state = kernel_split_state.ray_state;
int inactive_ray = dequeue_ray_index(QUEUE_INACTIVE_RAYS,
kernel_split_state.queue_data, kernel_split_params.queue_size, kernel_split_params.queue_index);
if(!IS_STATE(ray_state, inactive_ray, RAY_INACTIVE)) {
return false;
}
#define SPLIT_DATA_ENTRY(type, name, num) \
kernel_split_state.name[inactive_ray] = kernel_split_state.name[ray_index];
SPLIT_DATA_ENTRIES_BRANCHED_SHARED
#undef SPLIT_DATA_ENTRY
kernel_split_state.branched_state[inactive_ray].shared_sample_count = 0;
kernel_split_state.branched_state[inactive_ray].original_ray = ray_index;
kernel_split_state.branched_state[inactive_ray].waiting_on_shared_samples = false;
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
PathRadiance *inactive_L = &kernel_split_state.path_radiance[inactive_ray];
path_radiance_init(inactive_L, kernel_data.film.use_light_pass);
inactive_L->direct_throughput = L->direct_throughput;
path_radiance_copy_indirect(inactive_L, L);
ray_state[inactive_ray] = RAY_REGENERATED;
ADD_RAY_FLAG(ray_state, inactive_ray, RAY_BRANCHED_INDIRECT_SHARED);
ADD_RAY_FLAG(ray_state, inactive_ray, IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT));
atomic_fetch_and_inc_uint32((ccl_global uint*)&kernel_split_state.branched_state[ray_index].shared_sample_count);
return true;
}
/* bounce off surface and integrate indirect light */
ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter(KernelGlobals *kg,
int ray_index,
float num_samples_adjust,
ShaderData *saved_sd,
bool reset_path_state)
bool reset_path_state,
bool wait_for_shared)
{
SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
@ -155,12 +192,25 @@ ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter(
/* start the indirect path */
*tp *= num_samples_inv;
if(kernel_split_branched_indirect_start_shared(kg, ray_index)) {
continue;
}
return true;
}
branched_state->next_sample = 0;
}
branched_state->next_closure = sd->num_closure;
if(wait_for_shared) {
branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
if(branched_state->waiting_on_shared_samples) {
return true;
}
}
return false;
}

View File

@ -75,11 +75,30 @@ ccl_device_noinline bool kernel_split_branched_path_volume_indirect_light_iter(K
branched_state->next_sample = j+1;
branched_state->num_samples = num_samples;
/* Attempting to share too many samples is slow for volumes as it causes us to
* loop here more and have many calls to kernel_volume_integrate which evaluates
* shaders. The many expensive shader evaluations cause the work load to become
* unbalanced and many threads to become idle in this kernel. Limiting the
* number of shared samples here helps quite a lot.
*/
if(branched_state->shared_sample_count < 2) {
if(kernel_split_branched_indirect_start_shared(kg, ray_index)) {
continue;
}
}
return true;
}
# endif
}
branched_state->next_sample = num_samples;
branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
if(branched_state->waiting_on_shared_samples) {
return true;
}
kernel_split_branched_path_indirect_loop_end(kg, ray_index);
/* todo: avoid this calculation using decoupled ray marching */

View File

@ -0,0 +1,46 @@
/*
* Copyright 2011-2017 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
CCL_NAMESPACE_BEGIN
ccl_device void kernel_enqueue_inactive(KernelGlobals *kg,
ccl_local_param unsigned int *local_queue_atomics)
{
#ifdef __BRANCHED_PATH__
/* Enqeueue RAY_INACTIVE rays into QUEUE_INACTIVE_RAYS queue. */
if(ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
*local_queue_atomics = 0;
}
ccl_barrier(CCL_LOCAL_MEM_FENCE);
int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
char enqueue_flag = 0;
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_INACTIVE)) {
enqueue_flag = 1;
}
enqueue_ray_index_local(ray_index,
QUEUE_INACTIVE_RAYS,
enqueue_flag,
kernel_split_params.queue_size,
local_queue_atomics,
kernel_split_state.queue_data,
kernel_split_params.queue_index);
#endif /* __BRANCHED_PATH__ */
}
CCL_NAMESPACE_END

View File

@ -147,6 +147,7 @@ ccl_device void kernel_next_iteration_setup(KernelGlobals *kg,
ray_index,
1.0f,
&kernel_split_state.branched_state[ray_index].sd,
true,
true))
{
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED);
@ -193,6 +194,7 @@ ccl_device void kernel_next_iteration_setup(KernelGlobals *kg,
ray_index,
1.0f,
&kernel_split_state.branched_state[ray_index].sd,
true,
true))
{
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED);

View File

@ -43,11 +43,21 @@ ccl_device void kernel_scene_intersect(KernelGlobals *kg)
}
/* All regenerated rays become active here */
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED))
ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE);
if(IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) {
#ifdef __BRANCHED_PATH__
if(kernel_split_state.branched_state[ray_index].waiting_on_shared_samples) {
kernel_split_path_end(kg, ray_index);
}
else
#endif /* __BRANCHED_PATH__ */
{
ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE);
}
}
if(!IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE))
if(!IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) {
return;
}
#ifdef __KERNEL_DEBUG__
DebugData *debug_data = &kernel_split_state.debug_data[ray_index];

View File

@ -29,6 +29,14 @@ ccl_device void kernel_shadow_blocked_dl(KernelGlobals *kg)
kernel_split_state.queue_data, kernel_split_params.queue_size, 1);
}
#ifdef __BRANCHED_PATH__
/* TODO(mai): move this somewhere else? */
if(thread_index == 0) {
/* Clear QUEUE_INACTIVE_RAYS before next kernel. */
kernel_split_params.queue_index[QUEUE_INACTIVE_RAYS] = 0;
}
#endif /* __BRANCHED_PATH__ */
if(ray_index == QUEUE_EMPTY_SLOT)
return;

View File

@ -56,7 +56,20 @@ ccl_device_inline void kernel_split_path_end(KernelGlobals *kg, int ray_index)
ccl_global char *ray_state = kernel_split_state.ray_state;
#ifdef __BRANCHED_PATH__
if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT)) {
if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT_SHARED)) {
int orig_ray = kernel_split_state.branched_state[ray_index].original_ray;
PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
PathRadiance *orig_ray_L = &kernel_split_state.path_radiance[orig_ray];
path_radiance_sum_indirect(L);
path_radiance_accum_sample(orig_ray_L, L, 1);
atomic_fetch_and_dec_uint32((ccl_global uint*)&kernel_split_state.branched_state[orig_ray].shared_sample_count);
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_INACTIVE);
}
else if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT)) {
ASSIGN_RAY_STATE(ray_state, ray_index, RAY_LIGHT_INDIRECT_NEXT_ITER);
}
else if(IS_FLAG(ray_state, ray_index, RAY_BRANCHED_VOLUME_INDIRECT)) {

View File

@ -95,6 +95,10 @@ typedef ccl_global struct SplitBranchedState {
VolumeStack volume_stack[VOLUME_STACK_SIZE];
# endif /* __VOLUME__ */
#endif /*__SUBSURFACE__ */
int shared_sample_count; /* number of branched samples shared with other threads */
int original_ray; /* index of original ray when sharing branched samples */
bool waiting_on_shared_samples;
} SplitBranchedState;
#define SPLIT_DATA_BRANCHED_ENTRIES \
@ -137,6 +141,25 @@ typedef ccl_global struct SplitBranchedState {
SPLIT_DATA_BRANCHED_ENTRIES \
SPLIT_DATA_DEBUG_ENTRIES \
/* entries to be copied to inactive rays when sharing branched samples (TODO: which are actually needed?) */
#define SPLIT_DATA_ENTRIES_BRANCHED_SHARED \
SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \
SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \
SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \
SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \
SPLIT_DATA_ENTRY(ccl_global Intersection, isect, 1) \
SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \
SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 1) \
SPLIT_DATA_SUBSURFACE_ENTRIES \
SPLIT_DATA_VOLUME_ENTRIES \
SPLIT_DATA_BRANCHED_ENTRIES \
SPLIT_DATA_DEBUG_ENTRIES \
/* struct that holds pointers to data in the shared state buffer */
typedef struct SplitData {
#define SPLIT_DATA_ENTRY(type, name, num) type *name;

View File

@ -169,6 +169,7 @@ ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_it
ray_index,
num_samples_inv,
bssrdf_sd,
false,
false))
{
branched_state->ss_next_closure = i;
@ -187,6 +188,13 @@ ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_it
branched_state->ss_next_sample = 0;
}
branched_state->ss_next_closure = sd->num_closure;
branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
if(branched_state->waiting_on_shared_samples) {
return true;
}
kernel_split_branched_path_indirect_loop_end(kg, ray_index);
return false;

View File

@ -722,6 +722,7 @@ DeviceRequestedFeatures Session::get_requested_device_features()
requested_features.use_baking = bake_manager->get_baking();
requested_features.use_integrator_branched = (scene->integrator->method == Integrator::BRANCHED_PATH);
requested_features.use_transparent &= scene->integrator->transparent_shadows;
requested_features.use_denoising = params.use_denoising;
return requested_features;
}

View File

@ -35,6 +35,7 @@ ATOMIC_INLINE void atomic_update_max_z(size_t *maximum_value, size_t value)
#define atomic_add_and_fetch_float(p, x) atomic_add_and_fetch_fl((p), (x))
#define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1)
#define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_add_uint32((p), -1)
#define CCL_LOCAL_MEM_FENCE 0
#define ccl_barrier(flags) (void)0
@ -68,6 +69,7 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so
#define atomic_fetch_and_add_uint32(p, x) atomic_add((p), (x))
#define atomic_fetch_and_inc_uint32(p) atomic_inc((p))
#define atomic_fetch_and_dec_uint32(p) atomic_dec((p))
#define CCL_LOCAL_MEM_FENCE CLK_LOCAL_MEM_FENCE
#define ccl_barrier(flags) barrier(flags)
@ -79,7 +81,9 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so
#define atomic_add_and_fetch_float(p, x) (atomicAdd((float*)(p), (float)(x)) + (float)(x))
#define atomic_fetch_and_add_uint32(p, x) atomicAdd((unsigned int*)(p), (unsigned int)(x))
#define atomic_fetch_and_sub_uint32(p, x) atomicSub((unsigned int*)(p), (unsigned int)(x))
#define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1)
#define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_sub_uint32((p), 1)
#define CCL_LOCAL_MEM_FENCE
#define ccl_barrier(flags) __syncthreads()

View File

@ -31358,6 +31358,136 @@
fx="-0.78262758"
fy="294.63174"
r="6.6750002" />
<radialGradient
inkscape:collect="always"
xlink:href="#linearGradient33897"
id="radialGradient29130"
gradientUnits="userSpaceOnUse"
gradientTransform="matrix(1.6249996,2.7477764e-7,-3.1704883e-7,1.874986,-24.234082,-761.21063)"
cx="39.528847"
cy="871.2453"
fx="39.528847"
fy="871.2453"
r="2.0000005" />
<linearGradient
inkscape:collect="always"
id="linearGradient33897">
<stop
style="stop-color:#ffffff;stop-opacity:1;"
offset="0"
id="stop33893" />
<stop
style="stop-color:#ffffff;stop-opacity:0;"
offset="1"
id="stop33895" />
</linearGradient>
<linearGradient
inkscape:collect="always"
xlink:href="#linearGradient319-367"
id="linearGradient63713"
gradientUnits="userSpaceOnUse"
gradientTransform="matrix(0.8167109,0,0,0.8433415,239.34332,-149.78578)"
x1="104.90227"
y1="53.227627"
x2="114.94328"
y2="60.73848" />
<linearGradient
id="linearGradient319-367">
<stop
style="stop-color:#ffffff;stop-opacity:1;"
offset="0"
id="stop320-53" />
<stop
style="stop-color:#ffffff;stop-opacity:0;"
offset="1"
id="stop321-562" />
</linearGradient>
<linearGradient
inkscape:collect="always"
xlink:href="#linearGradient319-367"
id="linearGradient63715"
gradientUnits="userSpaceOnUse"
gradientTransform="translate(207,-246.99988)"
x1="-56.65625"
y1="342.03125"
x2="-53.1875"
y2="342.0625" />
<linearGradient
id="linearGradient3043">
<stop
style="stop-color:#ffffff;stop-opacity:1;"
offset="0"
id="stop3045" />
<stop
style="stop-color:#ffffff;stop-opacity:0;"
offset="1"
id="stop3047" />
</linearGradient>
<linearGradient
inkscape:collect="always"
xlink:href="#linearGradient319-367"
id="linearGradient63886"
x1="149.53125"
y1="95.781372"
x2="149.40625"
y2="103.12512"
gradientUnits="userSpaceOnUse" />
<linearGradient
id="linearGradient3050">
<stop
style="stop-color:#ffffff;stop-opacity:1;"
offset="0"
id="stop3052" />
<stop
style="stop-color:#ffffff;stop-opacity:0;"
offset="1"
id="stop3054" />
</linearGradient>
<linearGradient
inkscape:collect="always"
xlink:href="#linearGradient63892"
id="linearGradient63894"
x1="127.85783"
y1="115.03898"
x2="137.88899"
y2="121.44501"
gradientUnits="userSpaceOnUse" />
<linearGradient
inkscape:collect="always"
id="linearGradient63892">
<stop
style="stop-color:#ba5d00;stop-opacity:1"
offset="0"
id="stop63888" />
<stop
id="stop63896"
offset="0.5"
style="stop-color:#fa9a3a;stop-opacity:1" />
<stop
style="stop-color:#ffffff;stop-opacity:1"
offset="1"
id="stop63890" />
</linearGradient>
<linearGradient
inkscape:collect="always"
xlink:href="#linearGradient319-367"
id="linearGradient63719"
gradientUnits="userSpaceOnUse"
x1="132"
y1="117.26753"
x2="142.72656"
y2="127.72736" />
<linearGradient
id="linearGradient3062">
<stop
style="stop-color:#ffffff;stop-opacity:1;"
offset="0"
id="stop3064" />
<stop
style="stop-color:#ffffff;stop-opacity:0;"
offset="1"
id="stop3066" />
</linearGradient>
</defs>
<sodipodi:namedview
id="base"
@ -92721,6 +92851,406 @@
style="display:inline;overflow:visible;visibility:visible;fill:none;stroke:#a8df84;stroke-width:2.93474906;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;enable-background:accumulate"
id="path39834-2"
inkscape:connector-curvature="0" />
<g
id="g1418"
transform="matrix(1.6674711,0,0,1.6674711,-416.35106,-776.00982)"
style="display:inline;enable-background:new">
<path
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1374"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="4.2499995"
sodipodi:ry="4.2499995"
sodipodi:start="6.1086524"
sodipodi:end="7.5049158"
sodipodi:open="true"
d="m 45.185432,872.62412 c 0.358149,2.03116 -0.793733,4.02628 -2.731847,4.73169" />
<path
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1376"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="7"
sodipodi:ry="7"
sodipodi:start="4.4505896"
sodipodi:end="4.9741884"
sodipodi:open="true"
d="m 39.188267,866.60064 c 1.186888,-0.31802 2.436579,-0.31802 3.623467,0" />
<path
d="m 47.76148,875.17385 c -0.318025,1.18689 -0.942871,2.26916 -1.811733,3.13802"
sodipodi:open="true"
sodipodi:end="0.78539816"
sodipodi:start="0.26179939"
sodipodi:ry="6.9999995"
sodipodi:rx="6.9999995"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1378"
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
d="m 39.546414,877.35581 c -1.938113,-0.70541 -3.089995,-2.70053 -2.731846,-4.73169"
sodipodi:open="true"
sodipodi:end="3.3161256"
sodipodi:start="1.9198622"
sodipodi:ry="4.2499995"
sodipodi:rx="4.2499995"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1380"
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1382"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="6.9999995"
sodipodi:ry="6.9999995"
sodipodi:start="2.3561945"
sodipodi:end="2.8797933"
sodipodi:open="true"
d="m 36.050253,878.31187 c -0.868862,-0.86886 -1.493708,-1.95113 -1.811733,-3.13802" />
<rect
ry="2.75"
rx="2.75"
y="870.61212"
x="38.25"
height="5.5"
width="5.5000005"
id="rect1384"
style="fill:#4b2f1e;fill-opacity:1;fill-rule:nonzero;stroke:none" />
<rect
style="fill:#000000;fill-opacity:0;fill-rule:nonzero;stroke:none"
id="rect1386"
width="20"
height="20"
x="31"
y="862.36212"
rx="0"
ry="0" />
<rect
style="fill:#ff7d1f;fill-opacity:1;fill-rule:nonzero;stroke:none"
id="rect1388"
width="4.0000005"
height="4"
x="39"
y="871.36212"
rx="2"
ry="2" />
<rect
ry="1.5"
rx="1.5000004"
y="871.86212"
x="39.5"
height="3"
width="3.0000007"
id="rect1390"
style="fill:none;stroke:url(#radialGradient29130);stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1392"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="4.2499995"
sodipodi:ry="4.2499995"
sodipodi:start="1.9198622"
sodipodi:end="3.3161256"
sodipodi:open="true"
d="m 39.546414,877.35581 c -1.938113,-0.70541 -3.089995,-2.70053 -2.731846,-4.73169" />
<path
d="m 36.050253,878.31187 c -0.868862,-0.86886 -1.493708,-1.95113 -1.811733,-3.13802"
sodipodi:open="true"
sodipodi:end="2.8797933"
sodipodi:start="2.3561945"
sodipodi:ry="6.9999995"
sodipodi:rx="6.9999995"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1394"
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
d="m 38.875,877.04273 c -0.321653,-0.18571 -0.617575,-0.41278 -0.880204,-0.6754"
sodipodi:open="true"
sodipodi:end="2.3561945"
sodipodi:start="2.0943951"
sodipodi:ry="4.2499995"
sodipodi:rx="4.2499995"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1396"
style="opacity:0.7;fill:none;stroke:#ff8919;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
style="opacity:0.7;fill:none;stroke:#ff8919;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1398"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="6.9999995"
sodipodi:ry="6.9999995"
sodipodi:start="2.3561945"
sodipodi:end="2.6179939"
sodipodi:open="true"
d="m 36.050253,878.31187 c -0.432565,-0.43257 -0.806561,-0.91997 -1.11243,-1.44975" />
<path
style="fill:none;stroke:#4b2f1e;stroke-width:2.75;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1400"
sodipodi:type="arc"
sodipodi:cx="-41"
sodipodi:cy="-873.36212"
sodipodi:rx="4.2500005"
sodipodi:ry="4.2500005"
sodipodi:start="0.87266463"
sodipodi:end="2.268928"
sodipodi:open="true"
d="m -38.268152,-870.10643 c -1.579966,1.32575 -3.88373,1.32575 -5.463696,0"
inkscape:transform-center-x="-2.2499995"
inkscape:transform-center-y="-2.25"
transform="scale(-1,-1)" />
<path
d="m -38.268152,-870.10643 c -1.579966,1.32575 -3.88373,1.32575 -5.463696,0"
sodipodi:open="true"
sodipodi:end="2.268928"
sodipodi:start="0.87266463"
sodipodi:ry="4.2500005"
sodipodi:rx="4.2500005"
sodipodi:cy="-873.36212"
sodipodi:cx="-41"
sodipodi:type="arc"
id="path1402"
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
inkscape:transform-center-x="-2.2499995"
inkscape:transform-center-y="-2.25"
transform="scale(-1,-1)" />
<path
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1404"
sodipodi:type="arc"
sodipodi:cx="-41"
sodipodi:cy="-873.36212"
sodipodi:rx="7.0000005"
sodipodi:ry="7.0000005"
sodipodi:start="1.3089969"
sodipodi:end="1.8325957"
sodipodi:open="true"
d="m -39.188266,-866.60064 c -1.186888,0.31803 -2.436579,0.31803 -3.623467,0"
inkscape:transform-center-y="-3.7500001"
transform="scale(-1,-1)"
inkscape:transform-center-x="-3.7499995" />
<path
style="opacity:0.3;fill:none;stroke:#ffe4cb;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1406"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="4.25"
sodipodi:ry="4.25"
sodipodi:start="4.0142573"
sodipodi:end="4.9741884"
sodipodi:open="true"
d="m 38.268153,870.10643 c 1.062216,-0.8913 2.492452,-1.20838 3.831828,-0.84949" />
<path
d="m 39.188267,866.60064 c 1.186888,-0.31802 2.436579,-0.31802 3.623467,0"
sodipodi:open="true"
sodipodi:end="4.9741884"
sodipodi:start="4.4505896"
sodipodi:ry="7"
sodipodi:rx="7"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1408"
style="opacity:0.3;fill:none;stroke:#ffe4cb;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
style="opacity:0.6;fill:none;stroke:#ffffff;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1410"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="6.9999995"
sodipodi:ry="6.9999995"
sodipodi:start="4.4505896"
sodipodi:end="4.712389"
sodipodi:open="true"
d="M 39.188267,866.60064 C 39.779161,866.44231 40.388261,866.36212 41,866.36212" />
<path
d="m 38.268153,870.10643 c 0.475408,-0.39891 1.032412,-0.68887 1.631866,-0.84949"
sodipodi:open="true"
sodipodi:end="4.4505896"
sodipodi:start="4.0142573"
sodipodi:ry="4.25"
sodipodi:rx="4.25"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1412"
style="opacity:0.6;fill:none;stroke:#ffffff;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0" />
<path
sodipodi:open="true"
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="path1414"
sodipodi:type="arc"
sodipodi:cx="41"
sodipodi:cy="873.36212"
sodipodi:rx="7"
sodipodi:ry="7"
sodipodi:start="0.26179939"
sodipodi:end="0.78539816"
d="m 47.761481,875.17385 c -0.318026,1.18689 -0.942871,2.26916 -1.811734,3.13802" />
<path
d="m 45.185432,872.62412 c 0.358149,2.03116 -0.793733,4.02628 -2.731847,4.73169"
sodipodi:end="7.5049158"
sodipodi:start="6.1086524"
sodipodi:ry="4.2499995"
sodipodi:rx="4.2499995"
sodipodi:cy="873.36212"
sodipodi:cx="41"
sodipodi:type="arc"
id="path1416"
style="fill:none;stroke:#ffb36b;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
sodipodi:open="true" />
</g>
<g
transform="matrix(1.6674711,0,0,1.6674711,-138.14039,-258.06137)"
style="display:inline;enable-background:new"
id="g63699">
<path
inkscape:connector-curvature="0"
style="opacity:0.3;fill:none;stroke:#422200;stroke-width:0.80000001;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
d="m -102.5625,554.75 c -0.0429,0.005 -0.0849,0.0154 -0.125,0.0312 l -4.5,1.75 c -0.1919,0.0756 -0.31653,0.26258 -0.3125,0.4688 v 6.5 c 0.003,0.18741 0.11203,0.35691 0.28125,0.4375 l 4.5,2.25 c 0.13787,0.0682 0.29963,0.0682 0.4375,0 l 4.499997,-2.25 c 0.169224,-0.0806 0.278188,-0.25009 0.28125,-0.4375 V 557 c 0.004,-0.2062 -0.120622,-0.39315 -0.3125,-0.46875 l -4.499997,-1.75 c -0.0792,-0.0318 -0.16537,-0.0426 -0.25,-0.0312 z"
id="path63697" />
<g
id="g63687"
transform="translate(-252,462.99988)">
<g
transform="translate(-179,199.50012)"
id="g63679">
<path
inkscape:connector-curvature="0"
id="path63667"
d="m 328.5,-107.25 -4.5,1.75 v 6.5 l 4.5,2.25 4.5,-2.25 v -6.5 z"
style="fill:#422200;fill-opacity:1;fill-rule:evenodd;stroke:#422200;stroke-width:1.79999995;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
sodipodi:nodetypes="ccccccc" />
<g
transform="translate(179,-179)"
id="g63671">
<path
inkscape:connector-curvature="0"
inkscape:export-ydpi="90"
inkscape:export-xdpi="90"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png"
style="fill:#915515;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
d="m 154,80 v -6.5 l -4.5,-1.75 v 10.5 z"
id="path63669"
sodipodi:nodetypes="ccccc" />
</g>
<path
inkscape:connector-curvature="0"
sodipodi:nodetypes="ccccccc"
style="fill:#efa351;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
d="M 324,-99.00012 V -105.5 l 4.5,-1.75 0.5,0.25 v 10 l -0.5,0.25 z"
id="path63673"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png"
inkscape:export-xdpi="90"
inkscape:export-ydpi="90" />
<path
inkscape:connector-curvature="0"
style="fill:none;stroke:url(#linearGradient63713);stroke-width:0.99999982px;stroke-linecap:butt;stroke-linejoin:miter;stroke-opacity:1"
d="m 332.5,-105.5 v 6.25 l -4,2 -4,-2.00012 V -105.5"
id="path63675"
sodipodi:nodetypes="ccccc" />
<path
inkscape:connector-curvature="0"
inkscape:export-ydpi="90"
inkscape:export-xdpi="90"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\4.png"
sodipodi:nodetypes="ccccc"
style="fill:#f5ca9b;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
d="m 324,-105.5 4.5,-1.75 4.5,1.75 -4.5,2 z"
id="path63677" />
</g>
<path
inkscape:connector-curvature="0"
sodipodi:nodetypes="ccc"
style="opacity:0.95999995;fill:none;stroke:url(#linearGradient63715);stroke-width:1px;stroke-linecap:round;stroke-linejoin:round;stroke-opacity:1"
d="m 145.5,94.25012 c 0,0 4,1.75 4,1.75 l 4,-1.75"
id="path63681" />
<rect
style="opacity:0.25;fill:url(#linearGradient63886);fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1.06666696px;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
id="rect63683"
width="1"
height="6.7500019"
x="149"
y="96.000122" />
</g>
<rect
style="opacity:0;fill:#b3b3b3;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:2.79999995;marker:none;visibility:visible;display:inline;overflow:visible;enable-background:accumulate"
id="rect63665"
width="16"
height="16"
x="-113"
y="554" />
<g
style="display:inline;enable-background:new"
id="g63695"
transform="matrix(0.7882544,0,0,0.7883038,-210.45268,388.9974)"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\BLENDER ICONSET\blender-cvs-windows\.blender\.blender\icons\jendrzych's iconset.png"
inkscape:export-xdpi="90"
inkscape:export-ydpi="90">
<circle
sodipodi:ry="8"
sodipodi:rx="8"
sodipodi:cy="118"
sodipodi:cx="132"
r="8"
cy="118"
cx="132"
inkscape:export-ydpi="90"
inkscape:export-xdpi="90"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png"
id="circle63689"
style="fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:#422200;stroke-width:1.97436094;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
transform="matrix(0.6425292,0,0,0.642531,44.523834,146.81699)"
d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" />
<circle
sodipodi:ry="8"
sodipodi:rx="8"
sodipodi:cy="118"
sodipodi:cx="132"
r="8"
cy="118"
cx="132"
style="opacity:0.8;fill:url(#linearGradient63894);fill-opacity:1;fill-rule:nonzero;stroke:none;display:inline"
id="circle63691"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png"
inkscape:export-xdpi="90"
inkscape:export-ydpi="90"
transform="matrix(-0.5858806,-0.06590218,0.06677852,-0.5812167,198.80048,299.96262)"
d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" />
<circle
sodipodi:ry="8"
sodipodi:rx="8"
sodipodi:cy="118"
sodipodi:cx="132"
r="8"
cy="118"
cx="132"
transform="matrix(0.4991181,0,0,0.4991107,63.460522,163.7471)"
style="opacity:0.5;fill:none;stroke:url(#linearGradient63719);stroke-width:2.54167628;stroke-linecap:square;stroke-linejoin:round;stroke-miterlimit:4;stroke-opacity:1;stroke-dasharray:none;stroke-dashoffset:0"
id="circle63693"
inkscape:export-filename="C:\Documents and Settings\Tata\Pulpit\Kopia blender\.blender\icons\blender's iconset.png"
inkscape:export-xdpi="90"
inkscape:export-ydpi="90"
d="m 140,118 c 0,4.41828 -3.58172,8 -8,8 -4.41828,0 -8,-3.58172 -8,-8 0,-4.41828 3.58172,-8 8,-8 4.41828,0 8,3.58172 8,8 z" />
</g>
</g>
</g>
<path
id="path39836-9"

Before

Width:  |  Height:  |  Size: 4.4 MiB

After

Width:  |  Height:  |  Size: 4.4 MiB

View File

@ -118,7 +118,7 @@ MSG_COMMENT_PREFIX = "#~ "
MSG_CONTEXT_PREFIX = "MSGCTXT:"
# The default comment prefix used in po's.
PO_COMMENT_PREFIX= "# "
PO_COMMENT_PREFIX = "# "
# The comment prefix used to mark sources of msgids, in po's.
PO_COMMENT_PREFIX_SOURCE = "#: "
@ -130,7 +130,7 @@ PO_COMMENT_PREFIX_SOURCE_CUSTOM = "#. :src: "
PO_COMMENT_PREFIX_GENERATED = "#. "
# The comment prefix used to comment entries in po's.
PO_COMMENT_PREFIX_MSG= "#~ "
PO_COMMENT_PREFIX_MSG = "#~ "
# The comment prefix used to mark fuzzy msgids, in po's.
PO_COMMENT_FUZZY = "#, fuzzy"

View File

@ -124,9 +124,10 @@ class MATERIAL_PT_context_material(MaterialButtonsPanel, Panel):
ob = context.object
slot = context.material_slot
space = context.space_data
is_sortable = (len(ob.material_slots) > 1)
if ob:
is_sortable = (len(ob.material_slots) > 1)
rows = 1
if is_sortable:
rows = 4

View File

@ -338,21 +338,21 @@ class SEQUENCER_MT_add(Menu):
if len(bpy.data.scenes) > 10:
layout.operator_context = 'INVOKE_DEFAULT'
layout.operator("sequencer.scene_strip_add", text="Scene")
layout.operator("sequencer.scene_strip_add", text="Scene...")
else:
layout.operator_menu_enum("sequencer.scene_strip_add", "scene", text="Scene...")
layout.operator_menu_enum("sequencer.scene_strip_add", "scene", text="Scene")
if len(bpy.data.movieclips) > 10:
layout.operator_context = 'INVOKE_DEFAULT'
layout.operator("sequencer.movieclip_strip_add", text="Clips")
layout.operator("sequencer.movieclip_strip_add", text="Clips...")
else:
layout.operator_menu_enum("sequencer.movieclip_strip_add", "clip", text="Clip...")
layout.operator_menu_enum("sequencer.movieclip_strip_add", "clip", text="Clip")
if len(bpy.data.masks) > 10:
layout.operator_context = 'INVOKE_DEFAULT'
layout.operator("sequencer.mask_strip_add", text="Masks")
layout.operator("sequencer.mask_strip_add", text="Masks...")
else:
layout.operator_menu_enum("sequencer.mask_strip_add", "mask", text="Mask...")
layout.operator_menu_enum("sequencer.mask_strip_add", "mask", text="Mask")
layout.operator("sequencer.movie_strip_add", text="Movie")
layout.operator("sequencer.image_strip_add", text="Image")

View File

@ -250,7 +250,7 @@ def marker_menu_generic(layout):
layout.operator_context = 'INVOKE_DEFAULT'
layout.operator("marker.make_links_scene", text="Duplicate Marker to Scene...", icon='OUTLINER_OB_EMPTY')
else:
layout.operator_menu_enum("marker.make_links_scene", "scene", text="Duplicate Marker to Scene...")
layout.operator_menu_enum("marker.make_links_scene", "scene", text="Duplicate Marker to Scene")
layout.operator("marker.delete", text="Delete Marker")

View File

@ -129,7 +129,7 @@ class USERPREF_MT_app_templates(Menu):
text=bpy.path.display_name(d),
)
props.use_splash = True
props.app_template = d;
props.app_template = d
if use_install:
layout.separator()
@ -561,7 +561,7 @@ class USERPREF_PT_system(Panel):
# 3. Column
column = split.column()
column.label(text="Solid OpenGL lights:")
column.label(text="Solid OpenGL Lights:")
split = column.split(percentage=0.1)
split.label()

View File

@ -1266,14 +1266,14 @@ class INFO_MT_add(Menu):
layout.menu("INFO_MT_lightprobe_add")
layout.separator()
layout.operator_menu_enum("object.effector_add", "type", text="Force Field", icon='OUTLINER_OB_EMPTY')
layout.operator_menu_enum("object.effector_add", "type", text="Force Field", icon='OUTLINER_OB_FORCE_FIELD')
layout.separator()
if len(bpy.data.groups) > 10:
layout.operator_context = 'INVOKE_REGION_WIN'
layout.operator("object.group_instance_add", text="Group Instance...", icon='OUTLINER_OB_EMPTY')
layout.operator("object.group_instance_add", text="Group Instance...", icon='OUTLINER_OB_GROUP_INSTANCE')
else:
layout.operator_menu_enum("object.group_instance_add", "group", text="Group Instance", icon='OUTLINER_OB_EMPTY')
layout.operator_menu_enum("object.group_instance_add", "group", text="Group Instance", icon='OUTLINER_OB_GROUP_INSTANCE')
# ********** Object menu **********
@ -1654,7 +1654,7 @@ class VIEW3D_MT_make_links(Menu):
layout.operator("object.make_links_scene", text="Objects to Scene...", icon='OUTLINER_OB_EMPTY')
else:
layout.operator_context = 'EXEC_REGION_WIN'
layout.operator_menu_enum("object.make_links_scene", "scene", text="Objects to Scene...")
layout.operator_menu_enum("object.make_links_scene", "scene", text="Objects to Scene")
layout.operator_context = operator_context_default
layout.operator_enum("object.make_links_data", "type") # inline

View File

@ -89,12 +89,6 @@ typedef enum eAnimData_MergeCopy_Modes {
void BKE_animdata_merge_copy(struct ID *dst_id, struct ID *src_id, eAnimData_MergeCopy_Modes action_mode, bool fix_drivers);
/* Make Local */
void BKE_animdata_make_local(struct AnimData *adt);
/* Re-Assign ID's */
void BKE_animdata_relink(struct AnimData *adt);
/* ************************************* */
/* KeyingSets API */

View File

@ -399,73 +399,6 @@ void BKE_animdata_merge_copy(ID *dst_id, ID *src_id, eAnimData_MergeCopy_Modes a
}
}
/* Make Local -------------------------------------------- */
static void make_local_strips(ListBase *strips)
{
NlaStrip *strip;
for (strip = strips->first; strip; strip = strip->next) {
if (strip->act) BKE_action_make_local(G.main, strip->act, false);
if (strip->remap && strip->remap->target) BKE_action_make_local(G.main, strip->remap->target, false);
make_local_strips(&strip->strips);
}
}
/* Use local copy instead of linked copy of various ID-blocks */
void BKE_animdata_make_local(AnimData *adt)
{
NlaTrack *nlt;
/* Actions - Active and Temp */
if (adt->action) BKE_action_make_local(G.main, adt->action, false);
if (adt->tmpact) BKE_action_make_local(G.main, adt->tmpact, false);
/* Remaps */
if (adt->remap && adt->remap->target) BKE_action_make_local(G.main, adt->remap->target, false);
/* Drivers */
/* TODO: need to remap the ID-targets too? */
/* NLA Data */
for (nlt = adt->nla_tracks.first; nlt; nlt = nlt->next)
make_local_strips(&nlt->strips);
}
/* When duplicating data (i.e. objects), drivers referring to the original data will
* get updated to point to the duplicated data (if drivers belong to the new data)
*/
void BKE_animdata_relink(AnimData *adt)
{
/* sanity check */
if (adt == NULL)
return;
/* drivers */
if (adt->drivers.first) {
FCurve *fcu;
/* check each driver against all the base paths to see if any should go */
for (fcu = adt->drivers.first; fcu; fcu = fcu->next) {
ChannelDriver *driver = fcu->driver;
DriverVar *dvar;
/* driver variables */
for (dvar = driver->variables.first; dvar; dvar = dvar->next) {
/* only change the used targets, since the others will need fixing manually anyway */
DRIVER_TARGETS_USED_LOOPER(dvar)
{
if (dtar->id && dtar->id->newid) {
dtar->id = dtar->id->newid;
}
}
DRIVER_TARGETS_LOOPER_END
}
}
}
}
/* Sub-ID Regrouping ------------------------------------------- */
/**

View File

@ -790,7 +790,7 @@ static void default_get_tarmat(bConstraint *con, bConstraintOb *UNUSED(cob), bCo
ct = ctn; \
} \
} (void)0
/* --------- ChildOf Constraint ------------ */
static void childof_new_data(void *cdata)

View File

@ -986,19 +986,19 @@ static void do_physical_effector(EffectorCache *eff, EffectorData *efd, Effected
*/
void pdDoEffectors(ListBase *effectors, ListBase *colliders, EffectorWeights *weights, EffectedPoint *point, float *force, float *impulse)
{
/*
* Modifies the force on a particle according to its
* relation with the effector object
* Different kind of effectors include:
* Forcefields: Gravity-like attractor
* (force power is related to the inverse of distance to the power of a falloff value)
* Vortex fields: swirling effectors
* (particles rotate around Z-axis of the object. otherwise, same relation as)
* (Forcefields, but this is not done through a force/acceleration)
* Guide: particles on a path
* (particles are guided along a curve bezier or old nurbs)
* (is independent of other effectors)
*/
/*
* Modifies the force on a particle according to its
* relation with the effector object
* Different kind of effectors include:
* Forcefields: Gravity-like attractor
* (force power is related to the inverse of distance to the power of a falloff value)
* Vortex fields: swirling effectors
* (particles rotate around Z-axis of the object. otherwise, same relation as)
* (Forcefields, but this is not done through a force/acceleration)
* Guide: particles on a path
* (particles are guided along a curve bezier or old nurbs)
* (is independent of other effectors)
*/
EffectorCache *eff;
EffectorData efd;
int p=0, tot = 1, step = 1;

View File

@ -2865,7 +2865,7 @@ bool BKE_image_is_stereo(Image *ima)
{
return BKE_image_is_multiview(ima) &&
(BLI_findstring(&ima->views, STEREO_LEFT_NAME, offsetof(ImageView, name)) &&
BLI_findstring(&ima->views, STEREO_RIGHT_NAME, offsetof(ImageView, name)));
BLI_findstring(&ima->views, STEREO_RIGHT_NAME, offsetof(ImageView, name)));
}
static void image_init_multilayer_multiview(Image *ima, RenderResult *rr)

View File

@ -334,10 +334,10 @@ void BKE_ocean_eval_uv(struct Ocean *oc, struct OceanResult *ocr, float u, float
i1 = i1 % oc->_M;
j1 = j1 % oc->_N;
#define BILERP(m) (interpf(interpf(m[i1 * oc->_N + j1], m[i0 * oc->_N + j1], frac_x), \
interpf(m[i1 * oc->_N + j0], m[i0 * oc->_N + j0], frac_x), \
frac_z))
{
if (oc->_do_disp_y) {
ocr->disp[1] = BILERP(oc->_disp_y);

View File

@ -592,7 +592,7 @@ void psys_free(Object *ob, ParticleSystem *psys)
BLI_bvhtree_free(psys->bvhtree);
BLI_kdtree_free(psys->tree);
if (psys->fluid_springs)
MEM_freeN(psys->fluid_springs);

View File

@ -1114,7 +1114,7 @@ static void pbvh_update_draw_buffers(PBVH *bvh, PBVHNode **nodes, int totnode)
GPU_pbvh_bmesh_buffers_build(bvh->flags & PBVH_DYNTOPO_SMOOTH_SHADING);
break;
}
node->flag &= ~PBVH_RebuildDrawBuffers;
}

View File

@ -684,7 +684,7 @@ static float invGammaCorrect(float c)
else if (i >= RE_GAMMA_TABLE_SIZE) res = powf(c, valid_inv_gamma);
else res = inv_gamma_range_table[i] +
((c - color_domain_table[i]) * inv_gamfactor_table[i]);
return res;
}

View File

@ -1777,7 +1777,7 @@ static ImBuf *seq_proxy_fetch(const SeqRenderData *context, Sequence *seq, int c
if (proxy->anim == NULL) {
return NULL;
}
seq_open_anim_file(context->scene, seq, true);
sanim = seq->anims.first;
@ -1785,7 +1785,7 @@ static ImBuf *seq_proxy_fetch(const SeqRenderData *context, Sequence *seq, int c
return IMB_anim_absolute(proxy->anim, frameno, IMB_TC_NONE, IMB_PROXY_NONE);
}
if (seq_proxy_get_fname(ed, seq, cfra, render_size, name, context->view_id) == 0) {
return NULL;
}

View File

@ -2235,9 +2235,9 @@ static void sb_cf_threads_run(Scene *scene, Object *ob, float forcetime, float t
static void softbody_calc_forcesEx(Scene *scene, SceneLayer *sl, Object *ob, float forcetime, float timenow)
{
/* rule we never alter free variables :bp->vec bp->pos in here !
* this will ruin adaptive stepsize AKA heun! (BM)
*/
/* rule we never alter free variables :bp->vec bp->pos in here !
* this will ruin adaptive stepsize AKA heun! (BM)
*/
SoftBody *sb= ob->soft; /* is supposed to be there */
/*BodyPoint *bproot;*/ /* UNUSED */
ListBase *do_effector = NULL;

View File

@ -203,7 +203,7 @@ static float get_animated_scaleinf(StabContext *ctx, int framenr)
static void get_animated_target_pos(StabContext *ctx,
int framenr,
float target_pos[2])
float target_pos[2])
{
target_pos[0] = fetch_from_fcurve(ctx->target_pos[0],
framenr,

View File

@ -137,8 +137,8 @@ void _bli_array_grow_func(void **arr_p, const void *arr_static,
#define BLI_array_free(arr) \
if (arr && (char *)arr != _##arr##_static) { \
BLI_array_fake_user(arr); \
MEM_freeN(arr); \
BLI_array_fake_user(arr); \
MEM_freeN(arr); \
} (void)0
#define BLI_array_pop(arr) ( \

View File

@ -36,7 +36,7 @@
*/
#ifdef __cplusplus
extern "C" {
extern "C" {
#endif
struct BVHTree;
@ -62,7 +62,7 @@ typedef struct BVHTreeNearest {
int index; /* the index of the nearest found (untouched if none is found within a dist radius from the given coordinates) */
float co[3]; /* nearest coordinates (untouched it none is found within a dist radius from the given coordinates) */
float no[3]; /* normal at nearest coordinates (untouched it none is found within a dist radius from the given coordinates) */
float dist_sq; /* squared distance to search arround */
float dist_sq; /* squared distance to search around */
int flags;
} BVHTreeNearest;

View File

@ -62,6 +62,13 @@
/* used for iterative_raycast */
// #define USE_SKIP_LINKS
/* Use to print balanced output. */
// #define USE_PRINT_TREE
/* Check tree is valid. */
// #define USE_VERIFY_TREE
#define MAX_TREETYPE 32
/* Setting zero so we can catch bugs in BLI_task/KDOPBVH.
@ -571,10 +578,12 @@ static void node_join(BVHTree *tree, BVHNode *node)
}
}
/*
#ifdef USE_PRINT_TREE
/**
* Debug and information functions
*/
#if 0
static void bvhtree_print_tree(BVHTree *tree, BVHNode *node, int depth)
{
int i;
@ -597,30 +606,29 @@ static void bvhtree_print_tree(BVHTree *tree, BVHNode *node, int depth)
static void bvhtree_info(BVHTree *tree)
{
printf("BVHTree info\n");
printf("tree_type = %d, axis = %d, epsilon = %f\n",
printf("BVHTree Info: tree_type = %d, axis = %d, epsilon = %f\n",
tree->tree_type, tree->axis, tree->epsilon);
printf("nodes = %d, branches = %d, leafs = %d\n",
tree->totbranch + tree->totleaf, tree->totbranch, tree->totleaf);
printf("Memory per node = %ldbytes\n",
sizeof(BVHNode) + sizeof(BVHNode *) * tree->tree_type + sizeof(float) * tree->axis);
printf("BV memory = %dbytes\n",
(int)MEM_allocN_len(tree->nodebv));
printf("Memory per node = %ubytes\n",
(uint)(sizeof(BVHNode) + sizeof(BVHNode *) * tree->tree_type + sizeof(float) * tree->axis));
printf("BV memory = %ubytes\n",
(uint)MEM_allocN_len(tree->nodebv));
printf("Total memory = %ldbytes\n", sizeof(BVHTree) +
MEM_allocN_len(tree->nodes) +
MEM_allocN_len(tree->nodearray) +
MEM_allocN_len(tree->nodechild) +
MEM_allocN_len(tree->nodebv));
printf("Total memory = %ubytes\n",
(uint)(sizeof(BVHTree) +
MEM_allocN_len(tree->nodes) +
MEM_allocN_len(tree->nodearray) +
MEM_allocN_len(tree->nodechild) +
MEM_allocN_len(tree->nodebv)));
// bvhtree_print_tree(tree, tree->nodes[tree->totleaf], 0);
bvhtree_print_tree(tree, tree->nodes[tree->totleaf], 0);
}
#endif
#endif /* USE_PRINT_TREE */
#if 0
#ifdef USE_VERIFY_TREE
static void verify_tree(BVHTree *tree)
static void bvhtree_verify(BVHTree *tree)
{
int i, j, check = 0;
@ -661,7 +669,7 @@ static void verify_tree(BVHTree *tree)
printf("branches: %d, leafs: %d, total: %d\n",
tree->totbranch, tree->totleaf, tree->totbranch + tree->totleaf);
}
#endif
#endif /* USE_VERIFY_TREE */
/* Helper data and structures to build a min-leaf generalized implicit tree
* This code can be easily reduced
@ -907,16 +915,24 @@ static void non_recursive_bvh_div_nodes(
/* Loop tree levels (log N) loops */
for (i = 1, depth = 1; i <= num_branches; i = i * tree_type + tree_offset, depth++) {
const int first_of_next_level = i * tree_type + tree_offset;
const int end_j = min_ii(first_of_next_level, num_branches + 1); /* index of last branch on this level */
const int i_stop = min_ii(first_of_next_level, num_branches + 1); /* index of last branch on this level */
/* Loop all branches on this level */
cb_data.first_of_next_level = first_of_next_level;
cb_data.i = i;
cb_data.depth = depth;
BLI_task_parallel_range(
i, end_j, &cb_data, non_recursive_bvh_div_nodes_task_cb,
num_leafs > KDOPBVH_THREAD_LEAF_THRESHOLD);
if (true) {
BLI_task_parallel_range(
i, i_stop, &cb_data, non_recursive_bvh_div_nodes_task_cb,
num_leafs > KDOPBVH_THREAD_LEAF_THRESHOLD);
}
else {
/* Less hassle for debugging. */
for (int i_task = i; i_task < i_stop; i_task++) {
non_recursive_bvh_div_nodes_task_cb(&cb_data, i_task);
}
}
}
}
@ -1050,7 +1066,13 @@ void BLI_bvhtree_balance(BVHTree *tree)
build_skip_links(tree, tree->nodes[tree->totleaf], NULL, NULL);
#endif
/* bvhtree_info(tree); */
#ifdef USE_VERIFY_TREE
bvhtree_verify(tree);
#endif
#ifdef USE_PRINT_TREE
bvhtree_info(tree);
#endif
}
void BLI_bvhtree_insert(BVHTree *tree, int index, const float co[3], int numpoints)

View File

@ -1560,7 +1560,7 @@ BArrayState *BLI_array_store_state_add(
const void *data, const size_t data_len,
const BArrayState *state_reference)
{
/* ensure we're aligned to the stride */
/* ensure we're aligned to the stride */
BLI_assert((data_len % bs->info.chunk_stride) == 0);
#ifdef USE_PARANOID_CHECKS

View File

@ -422,13 +422,13 @@ void BLI_file_free_lines(LinkNode *lines)
bool BLI_file_older(const char *file1, const char *file2)
{
#ifdef WIN32
struct _stat st1, st2;
struct _stat st1, st2;
UTF16_ENCODE(file1);
UTF16_ENCODE(file2);
if (_wstat(file1_16, &st1)) return false;
if (_wstat(file2_16, &st2)) return false;
if (_wstat(file2_16, &st2)) return false;
UTF16_UN_ENCODE(file2);
UTF16_UN_ENCODE(file1);

View File

@ -1127,7 +1127,7 @@ static void task_parallel_range_ex(
atomic_fetch_and_add_uint32((uint32_t *)(&state.iter), 0);
if (use_userdata_chunk) {
userdata_chunk_array = MALLOCA(userdata_chunk_size * num_tasks);
userdata_chunk_array = MALLOCA(userdata_chunk_size * num_tasks);
}
for (i = 0; i < num_tasks; i++) {

View File

@ -10374,7 +10374,7 @@ void BLO_library_link_copypaste(Main *mainl, BlendHandle *bh)
static ID *link_named_part_ex(
Main *mainl, FileData *fd, const short idcode, const char *name, const short flag,
Scene *scene, SceneLayer *sl, const bool use_placeholders, const bool force_indirect)
Scene *scene, SceneLayer *sl, const bool use_placeholders, const bool force_indirect)
{
ID *id = link_named_part(mainl, fd, idcode, name, use_placeholders, force_indirect);

View File

@ -153,7 +153,8 @@ void DEG_add_object_cache_relation(struct DepsNodeHandle *handle,
eDepsObjectComponentType component,
const char *description);
/* TODO(sergey): Remove once all geometry update is granular. */
struct Depsgraph *DEG_get_graph_from_handle(struct DepsNodeHandle *handle);
void DEG_add_special_eval_flag(struct Depsgraph *graph, struct ID *id, short flag);
/* Utility functions for physics modifiers */

View File

@ -372,6 +372,11 @@ void DepsgraphRelationBuilder::add_forcefield_relations(const OperationKey &key,
pdEndEffectors(&effectors);
}
Depsgraph *DepsgraphRelationBuilder::getGraph()
{
return m_graph;
}
/* **** Functions to build relations between entities **** */
void DepsgraphRelationBuilder::begin_build(Main *bmain)

View File

@ -249,6 +249,8 @@ struct DepsgraphRelationBuilder
template <typename KeyType>
OperationDepsNode *find_operation_node(const KeyType &key);
Depsgraph *getGraph();
protected:
RootDepsNode *find_node(const RootKey &key) const;
TimeSourceDepsNode *find_node(const TimeSourceKey &key) const;

View File

@ -166,6 +166,13 @@ void DEG_add_bone_relation(DepsNodeHandle *handle,
description);
}
struct Depsgraph *DEG_get_graph_from_handle(struct DepsNodeHandle *handle)
{
DEG::DepsNodeHandle *deg_handle = get_handle(handle);
DEG::DepsgraphRelationBuilder *relation_builder = deg_handle->builder;
return reinterpret_cast<Depsgraph *>(relation_builder->getGraph());
}
void DEG_add_special_eval_flag(Depsgraph *graph, ID *id, short flag)
{
DEG::Depsgraph *deg_graph = reinterpret_cast<DEG::Depsgraph *>(graph);

View File

@ -452,7 +452,7 @@ void ANIM_editkeyframes_refresh(bAnimContext *ac)
ok |= KEYFRAME_OK_H2; \
} \
} (void)0
/* ------------------------ */
static short ok_bezier_frame(KeyframeEditData *ked, BezTriple *bezt)

View File

@ -313,9 +313,9 @@ DEF_ICON(OUTLINER_OB_ARMATURE)
DEF_ICON(OUTLINER_OB_FONT)
DEF_ICON(OUTLINER_OB_SURFACE)
DEF_ICON(OUTLINER_OB_SPEAKER)
DEF_ICON(OUTLINER_OB_FORCE_FIELD)
DEF_ICON(OUTLINER_OB_GROUP_INSTANCE)
#ifndef DEF_ICON_BLANK_SKIP
DEF_ICON(BLANK120)
DEF_ICON(BLANK121)
DEF_ICON(BLANK122)
DEF_ICON(BLANK123)
DEF_ICON(BLANK124)

View File

@ -383,8 +383,8 @@ void WM_OT_alembic_export(wmOperatorType *ot)
"Enable this to run the import in the background, disable to block Blender while importing");
/* This dummy prop is used to check whether we need to init the start and
* end frame values to that of the scene's, otherwise they are reset at
* every change, draw update. */
* end frame values to that of the scene's, otherwise they are reset at
* every change, draw update. */
RNA_def_boolean(ot->srna, "init_scene_frame_range", false, "", "");
}

View File

@ -2055,24 +2055,6 @@ void ED_object_single_users(Main *bmain, Scene *scene, const bool full, const bo
/******************************* Make Local ***********************************/
/* helper for below, ma was checked to be not NULL */
static void make_local_makelocalmaterial(Material *ma)
{
AnimData *adt;
int b;
id_make_local(G.main, &ma->id, false, false);
for (b = 0; b < MAX_MTEX; b++)
if (ma->mtex[b] && ma->mtex[b]->tex)
id_make_local(G.main, &ma->mtex[b]->tex->id, false, false);
adt = BKE_animdata_from_id(&ma->id);
if (adt) BKE_animdata_make_local(adt);
/* nodetree? XXX */
}
enum {
MAKE_LOCAL_SELECT_OB = 1,
MAKE_LOCAL_SELECT_OBDATA = 2,
@ -2159,125 +2141,145 @@ static bool make_local_all__instance_indirect_unused(Main *bmain, Scene *scene,
return changed;
}
static void make_local_animdata_tag_strips(ListBase *strips)
{
NlaStrip *strip;
for (strip = strips->first; strip; strip = strip->next) {
if (strip->act) {
strip->act->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
if (strip->remap && strip->remap->target) {
strip->remap->target->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
make_local_animdata_tag_strips(&strip->strips);
}
}
/* Tag all actions used by given animdata to be made local. */
static void make_local_animdata_tag(AnimData *adt)
{
if (adt) {
/* Actions - Active and Temp */
if (adt->action) {
adt->action->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
if (adt->tmpact) {
adt->tmpact->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
/* Remaps */
if (adt->remap && adt->remap->target) {
adt->remap->target->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
/* Drivers */
/* TODO: need to handle the ID-targets too? */
/* NLA Data */
for (NlaTrack *nlt = adt->nla_tracks.first; nlt; nlt = nlt->next) {
make_local_animdata_tag_strips(&nlt->strips);
}
}
}
static void make_local_material_tag(Material *ma)
{
if (ma) {
ma->id.tag &= ~LIB_TAG_PRE_EXISTING;
make_local_animdata_tag(BKE_animdata_from_id(&ma->id));
/* About nodetrees: root one is made local together with material, others we keep linked for now... */
for (int a = 0; a < MAX_MTEX; a++) {
if (ma->mtex[a] && ma->mtex[a]->tex) {
ma->mtex[a]->tex->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
}
}
}
static int make_local_exec(bContext *C, wmOperator *op)
{
Main *bmain = CTX_data_main(C);
Scene *scene = CTX_data_scene(C);
SceneLayer *sl = CTX_data_scene_layer(C);
SceneCollection *sc = CTX_data_scene_collection(C);
AnimData *adt;
ParticleSystem *psys;
Material *ma, ***matarar;
Lamp *la;
ID *id;
const int mode = RNA_enum_get(op->ptr, "type");
int a, b;
int a;
/* Note: we (ab)use LIB_TAG_PRE_EXISTING to cherry pick which ID to make local... */
if (mode == MAKE_LOCAL_ALL) {
SceneLayer *sl = CTX_data_scene_layer(C);
SceneCollection *sc = CTX_data_scene_collection(C);
BKE_main_id_tag_all(bmain, LIB_TAG_PRE_EXISTING, false);
/* de-select so the user can differentiate newly instanced from existing objects */
BKE_scene_base_deselect_all(scene);
if (make_local_all__instance_indirect_unused(bmain, scene, sl, sc)) {
BKE_report(op->reports, RPT_INFO,
"Orphan library objects added to the current scene to avoid loss");
}
BKE_library_make_local(bmain, NULL, NULL, false, false); /* NULL is all libs */
WM_event_add_notifier(C, NC_WINDOW, NULL);
return OPERATOR_FINISHED;
}
tag_localizable_objects(C, mode);
CTX_DATA_BEGIN (C, Object *, ob, selected_objects)
{
if ((ob->id.tag & LIB_TAG_DOIT) == 0) {
continue;
}
if (ob->id.lib)
id_make_local(bmain, &ob->id, false, false);
}
CTX_DATA_END;
/* maybe object pointers */
CTX_DATA_BEGIN (C, Object *, ob, selected_objects)
{
if (ob->id.lib == NULL) {
ID_NEW_REMAP(ob->parent);
BKE_report(op->reports, RPT_INFO, "Orphan library objects added to the current scene to avoid loss");
}
}
CTX_DATA_END;
else {
BKE_main_id_tag_all(bmain, LIB_TAG_PRE_EXISTING, true);
tag_localizable_objects(C, mode);
CTX_DATA_BEGIN (C, Object *, ob, selected_objects)
{
if ((ob->id.tag & LIB_TAG_DOIT) == 0) {
continue;
}
id = ob->data;
if (id && (ELEM(mode, MAKE_LOCAL_SELECT_OBDATA, MAKE_LOCAL_SELECT_OBDATA_MATERIAL))) {
id_make_local(bmain, id, false, false);
adt = BKE_animdata_from_id(id);
if (adt) BKE_animdata_make_local(adt);
/* tag indirect data direct */
matarar = give_matarar(ob);
if (matarar) {
for (a = 0; a < ob->totcol; a++) {
ma = (*matarar)[a];
if (ma)
id_lib_extern(&ma->id);
}
}
}
for (psys = ob->particlesystem.first; psys; psys = psys->next)
id_make_local(bmain, &psys->part->id, false, false);
adt = BKE_animdata_from_id(&ob->id);
if (adt) BKE_animdata_make_local(adt);
}
CTX_DATA_END;
if (mode == MAKE_LOCAL_SELECT_OBDATA_MATERIAL) {
CTX_DATA_BEGIN (C, Object *, ob, selected_objects)
{
if ((ob->id.tag & LIB_TAG_DOIT) == 0) {
continue;
}
if (ob->type == OB_LAMP) {
la = ob->data;
for (b = 0; b < MAX_MTEX; b++)
if (la->mtex[b] && la->mtex[b]->tex)
id_make_local(bmain, &la->mtex[b]->tex->id, false, false);
ob->id.tag &= ~LIB_TAG_PRE_EXISTING;
make_local_animdata_tag(BKE_animdata_from_id(&ob->id));
for (psys = ob->particlesystem.first; psys; psys = psys->next) {
psys->part->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
else {
if (mode == MAKE_LOCAL_SELECT_OBDATA_MATERIAL) {
for (a = 0; a < ob->totcol; a++) {
ma = ob->mat[a];
if (ma)
make_local_makelocalmaterial(ma);
if (ma) {
make_local_material_tag(ma);
}
}
matarar = (Material ***)give_matarar(ob);
if (matarar) {
for (a = 0; a < ob->totcol; a++) {
ma = (*matarar)[a];
if (ma)
make_local_makelocalmaterial(ma);
if (ma) {
make_local_material_tag(ma);
}
}
}
if (ob->type == OB_LAMP) {
BLI_assert(ob->data != NULL);
la = ob->data;
for (a = 0; a < MAX_MTEX; a++) {
if (la->mtex[a] && la->mtex[a]->tex) {
la->id.tag &= ~LIB_TAG_PRE_EXISTING;
}
}
}
}
if (ELEM(mode, MAKE_LOCAL_SELECT_OBDATA, MAKE_LOCAL_SELECT_OBDATA_MATERIAL) && ob->data != NULL) {
ID *ob_data = ob->data;
ob_data->tag &= ~LIB_TAG_PRE_EXISTING;
make_local_animdata_tag(BKE_animdata_from_id(ob_data));
}
}
CTX_DATA_END;
}
BKE_main_id_clear_newpoins(bmain);
WM_event_add_notifier(C, NC_WINDOW, NULL);
BKE_library_make_local(bmain, NULL, NULL, true, false); /* NULL is all libs */
WM_event_add_notifier(C, NC_WINDOW, NULL);
return OPERATOR_FINISHED;
}

View File

@ -281,10 +281,10 @@ void DPAINT_OT_output_toggle(wmOperatorType *ot)
/***************************** Image Sequence Baking ******************************/
typedef struct DynamicPaintBakeJob {
/* from wmJob */
void *owner;
short *stop, *do_update;
float *progress;
/* from wmJob */
void *owner;
short *stop, *do_update;
float *progress;
struct Main *bmain;
Scene *scene;
@ -300,13 +300,13 @@ typedef struct DynamicPaintBakeJob {
static void dpaint_bake_free(void *customdata)
{
DynamicPaintBakeJob *job = customdata;
MEM_freeN(job);
DynamicPaintBakeJob *job = customdata;
MEM_freeN(job);
}
static void dpaint_bake_endjob(void *customdata)
{
DynamicPaintBakeJob *job = customdata;
DynamicPaintBakeJob *job = customdata;
DynamicPaintCanvasSettings *canvas = job->canvas;
canvas->flags &= ~MOD_DPAINT_BAKING;
@ -314,7 +314,7 @@ static void dpaint_bake_endjob(void *customdata)
dynamicPaint_freeSurfaceData(job->surface);
G.is_rendering = false;
BKE_spacedata_draw_locks(false);
BKE_spacedata_draw_locks(false);
WM_set_locked_interface(G.main->wm.first, false);
@ -424,26 +424,26 @@ static void dynamicPaint_bakeImageSequence(DynamicPaintBakeJob *job)
static void dpaint_bake_startjob(void *customdata, short *stop, short *do_update, float *progress)
{
DynamicPaintBakeJob *job = customdata;
DynamicPaintBakeJob *job = customdata;
job->stop = stop;
job->do_update = do_update;
job->progress = progress;
job->stop = stop;
job->do_update = do_update;
job->progress = progress;
job->start = PIL_check_seconds_timer();
job->success = 1;
G.is_break = false; /* reset BKE_blender_test_break*/
G.is_break = false; /* reset BKE_blender_test_break*/
/* XXX annoying hack: needed to prevent data corruption when changing
* scene frame in separate threads
*/
G.is_rendering = true;
BKE_spacedata_draw_locks(true);
*/
G.is_rendering = true;
BKE_spacedata_draw_locks(true);
dynamicPaint_bakeImageSequence(job);
*do_update = true;
*stop = 0;
*do_update = true;
*stop = 0;
}
/*

View File

@ -100,45 +100,45 @@ static int ptcache_job_break(void *customdata)
static void ptcache_job_update(void *customdata, float progress, int *cancel)
{
PointCacheJob *job = customdata;
PointCacheJob *job = customdata;
if (ptcache_job_break(job)) {
*cancel = 1;
}
if (ptcache_job_break(job)) {
*cancel = 1;
}
*(job->do_update) = true;
*(job->progress) = progress;
*(job->do_update) = true;
*(job->progress) = progress;
}
static void ptcache_job_startjob(void *customdata, short *stop, short *do_update, float *progress)
{
PointCacheJob *job = customdata;
PointCacheJob *job = customdata;
job->stop = stop;
job->do_update = do_update;
job->progress = progress;
job->stop = stop;
job->do_update = do_update;
job->progress = progress;
G.is_break = false;
G.is_break = false;
/* XXX annoying hack: needed to prevent data corruption when changing
* scene frame in separate threads
*/
G.is_rendering = true;
BKE_spacedata_draw_locks(true);
/* XXX annoying hack: needed to prevent data corruption when changing
* scene frame in separate threads
*/
G.is_rendering = true;
BKE_spacedata_draw_locks(true);
BKE_ptcache_bake(job->baker);
*do_update = true;
*stop = 0;
*do_update = true;
*stop = 0;
}
static void ptcache_job_endjob(void *customdata)
{
PointCacheJob *job = customdata;
PointCacheJob *job = customdata;
Scene *scene = job->baker->scene;
G.is_rendering = false;
BKE_spacedata_draw_locks(false);
G.is_rendering = false;
BKE_spacedata_draw_locks(false);
WM_set_locked_interface(G.main->wm.first, false);

View File

@ -552,7 +552,7 @@ void nla_buttons_register(ARegionType *art)
pt = MEM_callocN(sizeof(PanelType), "spacetype nla panel modifiers");
strcpy(pt->idname, "NLA_PT_modifiers");
strcpy(pt->label, N_("Modifiers"));
strcpy(pt->category, "Modifiers");
strcpy(pt->category, "Modifiers");
strcpy(pt->translation_context, BLT_I18NCONTEXT_DEFAULT_BPYRNA);
pt->draw = nla_panel_modifiers;
pt->poll = nla_strip_eval_panel_poll;

View File

@ -683,7 +683,7 @@ static int sequencer_select_less_exec(bContext *C, wmOperator *UNUSED(op))
if (!select_more_less_seq__internal(scene, false, false))
return OPERATOR_CANCELLED;
WM_event_add_notifier(C, NC_SCENE | ND_SEQUENCER | NA_SELECTED, scene);
return OPERATOR_FINISHED;

View File

@ -975,7 +975,7 @@ static void recalcData_sequencer(TransInfo *t)
/* force recalculation of triangles during transformation */
static void recalcData_gpencil_strokes(TransInfo *t)
{
{
TransData *td = t->data;
for (int i = 0; i < t->total; i++, td++) {
bGPDstroke *gps = td->extra;

View File

@ -2859,7 +2859,7 @@ static PBool p_chart_symmetry_pins(PChart *chart, PEdge *outer, PVert **pin1, PV
PEdge *cure = NULL, *firste1 = NULL, *firste2 = NULL, *nextbe;
float maxlen = 0.0f, curlen = 0.0f, totlen = 0.0f, firstlen = 0.0f;
float len1, len2;
/* find longest series of verts split in the chart itself, these are
* marked during construction */
be = outer;

View File

@ -6450,7 +6450,7 @@ static int rna_function_parameter_parse(PointerRNA *ptr, PropertyRNA *prop, Prop
tid, fid, pid, RNA_struct_identifier(ptype), RNA_struct_identifier(srna));
return -1;
}
*((void **)dest) = *((void **)src);
break;

Some files were not shown because too many files have changed in this diff Show More