forked from bartvdbraak/blender
Merge branch 'master' into blender2.8
This commit is contained in:
commit
0f4f4d8754
@ -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')
|
||||
|
@ -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",
|
||||
|
@ -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)
|
||||
|
@ -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");
|
||||
|
@ -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());
|
||||
|
@ -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,
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
};
|
||||
|
@ -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);
|
||||
|
@ -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;
|
||||
|
@ -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.
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
|
@ -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)
|
||||
|
@ -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));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
@ -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,
|
||||
|
@ -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
|
||||
|
@ -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__
|
||||
|
@ -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)
|
||||
|
@ -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__
|
||||
|
@ -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))
|
||||
|
@ -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)
|
||||
|
@ -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)
|
||||
|
@ -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)
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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
|
||||
|
||||
|
@ -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"
|
||||
|
72
intern/cycles/kernel/kernels/opencl/kernel_split_function.h
Normal file
72
intern/cycles/kernel/kernels/opencl/kernel_split_function.h
Normal 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
|
||||
|
@ -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
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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 */
|
||||
|
46
intern/cycles/kernel/split/kernel_enqueue_inactive.h
Normal file
46
intern/cycles/kernel/split/kernel_enqueue_inactive.h
Normal 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
|
@ -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);
|
||||
|
@ -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];
|
||||
|
@ -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;
|
||||
|
||||
|
@ -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)) {
|
||||
|
@ -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;
|
||||
|
@ -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;
|
||||
|
@ -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;
|
||||
}
|
||||
|
@ -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()
|
||||
|
@ -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 |
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
@ -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"
|
||||
|
@ -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
|
||||
|
@ -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")
|
||||
|
@ -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")
|
||||
|
||||
|
@ -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()
|
||||
|
@ -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
|
||||
|
@ -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 */
|
||||
|
||||
|
@ -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 ------------------------------------------- */
|
||||
|
||||
/**
|
||||
|
@ -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)
|
||||
|
@ -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;
|
||||
|
@ -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)
|
||||
|
@ -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);
|
||||
|
@ -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);
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
@ -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;
|
||||
|
@ -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,
|
||||
|
@ -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) ( \
|
||||
|
@ -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;
|
||||
|
||||
|
@ -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)
|
||||
|
@ -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
|
||||
|
@ -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);
|
||||
|
@ -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++) {
|
||||
|
@ -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);
|
||||
|
||||
|
@ -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 */
|
||||
|
@ -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)
|
||||
|
@ -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;
|
||||
|
@ -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);
|
||||
|
@ -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)
|
||||
|
@ -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)
|
||||
|
@ -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, "", "");
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
|
@ -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;
|
||||
}
|
||||
|
||||
/*
|
||||
|
@ -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);
|
||||
|
||||
|
@ -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;
|
||||
|
@ -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;
|
||||
|
@ -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;
|
||||
|
@ -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;
|
||||
|
@ -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
Loading…
Reference in New Issue
Block a user