diff options
102 files changed, 1655 insertions, 1256 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 12f546dd158..6e6520cfdda 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1277,6 +1277,7 @@ if(CMAKE_COMPILER_IS_GNUCC) ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_NO_DIV_BY_ZERO -Wno-div-by-zero) ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_TYPE_LIMITS -Wtype-limits) ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_FORMAT_SIGN -Wformat-signedness) + ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_RESTRICT -Wrestrict) # gcc 4.2 gives annoying warnings on every file with this if(NOT "${CMAKE_C_COMPILER_VERSION}" VERSION_LESS "4.3") @@ -1318,6 +1319,7 @@ if(CMAKE_COMPILER_IS_GNUCC) ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_WRITE_STRINGS -Wwrite-strings) ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_UNDEF -Wundef) ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_FORMAT_SIGN -Wformat-signedness) + ADD_CHECK_CXX_COMPILER_FLAG(CXX_WARNINGS CXX_WARN_RESTRICT -Wrestrict) # gcc 4.2 gives annoying warnings on every file with this if(NOT "${CMAKE_C_COMPILER_VERSION}" VERSION_LESS "4.3") @@ -1354,7 +1356,7 @@ elseif(CMAKE_C_COMPILER_ID MATCHES "Clang") set(CMAKE_REQUIRED_FLAGS "-L${LIBDIR}/openmp/lib -liomp5") # these are only used for the checks endif() - # strange, clang complains these are not supported, but then yses them. + # strange, clang complains these are not supported, but then uses them. ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_ALL -Wall) ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_ERROR_IMPLICIT_FUNCTION_DECLARATION -Werror=implicit-function-declaration) ADD_CHECK_C_COMPILER_FLAG(C_WARNINGS C_WARN_ERROR_RETURN_TYPE -Werror=return-type) diff --git a/build_files/build_environment/install_deps.sh b/build_files/build_environment/install_deps.sh index 56756e68852..07ccf0cf3a6 100755 --- a/build_files/build_environment/install_deps.sh +++ b/build_files/build_environment/install_deps.sh @@ -1152,7 +1152,7 @@ compile_Numpy() { cd $_src - $_python/bin/python3 setup.py install --prefix=$_inst + $_python/bin/python3 setup.py install --old-and-unmanageable --prefix=$_inst if [ -d $_inst ]; then # Can't use _create_inst_shortcut here... diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index 6a1106328fb..19e3c0a9075 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -171,15 +171,15 @@ public: DeviceRequestedFeatures requested_features; - KernelFunctions<void(*)(KernelGlobals *, float *, unsigned int *, int, int, int, int, int)> path_trace_kernel; - KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_half_float_kernel; - KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_byte_kernel; - KernelFunctions<void(*)(KernelGlobals *, uint4 *, float4 *, float*, int, int, int, int, int)> shader_kernel; + KernelFunctions<void(*)(KernelGlobals *, float *, int, int, int, int, int)> path_trace_kernel; + KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_half_float_kernel; + KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_byte_kernel; + KernelFunctions<void(*)(KernelGlobals *, uint4 *, float4 *, int, int, int, int, int)> shader_kernel; - KernelFunctions<void(*)(int, TilesInfo*, int, int, float*, float*, float*, float*, float*, int*, int, int, bool)> filter_divide_shadow_kernel; - KernelFunctions<void(*)(int, TilesInfo*, int, int, int, int, float*, float*, int*, int, int, bool)> filter_get_feature_kernel; - KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)> filter_detect_outliers_kernel; - KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)> filter_combine_halves_kernel; + KernelFunctions<void(*)(int, TilesInfo*, int, int, float*, float*, float*, float*, float*, int*, int, int)> filter_divide_shadow_kernel; + KernelFunctions<void(*)(int, TilesInfo*, int, int, int, int, float*, float*, int*, int, int)> filter_get_feature_kernel; + KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)> filter_detect_outliers_kernel; + KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)> filter_combine_halves_kernel; KernelFunctions<void(*)(int, int, float*, float*, float*, int*, int, int, float, float)> filter_nlm_calc_difference_kernel; KernelFunctions<void(*)(float*, float*, int*, int, int)> filter_nlm_blur_kernel; @@ -192,7 +192,7 @@ public: KernelFunctions<void(*)(int, int, int, int, int, float*, int*, float*, float3*, int*, int)> filter_finalize_kernel; KernelFunctions<void(*)(KernelGlobals *, ccl_constant KernelData*, ccl_global void*, int, ccl_global char*, - ccl_global uint*, int, int, int, int, int, int, int, int, ccl_global int*, int, + int, int, int, int, int, int, int, int, ccl_global int*, int, ccl_global char*, ccl_global unsigned int*, unsigned int, ccl_global float*)> data_init_kernel; unordered_map<string, KernelFunctions<void(*)(KernelGlobals*, KernelData*)> > split_kernels; @@ -563,8 +563,7 @@ public: (float*) buffer_variance_ptr, &task->rect.x, task->render_buffer.pass_stride, - task->render_buffer.denoising_data_offset, - use_split_kernel); + task->render_buffer.denoising_data_offset); } } return true; @@ -587,8 +586,7 @@ public: (float*) variance_ptr, &task->rect.x, task->render_buffer.pass_stride, - task->render_buffer.denoising_data_offset, - use_split_kernel); + task->render_buffer.denoising_data_offset); } } return true; @@ -617,7 +615,6 @@ public: void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg) { float *render_buffer = (float*)tile.buffer; - uint *rng_state = (uint*)tile.rng_state; int start_sample = tile.start_sample; int end_sample = tile.start_sample + tile.num_samples; @@ -629,7 +626,7 @@ public: for(int y = tile.y; y < tile.y + tile.h; y++) { for(int x = tile.x; x < tile.x + tile.w; x++) { - path_trace_kernel()(kg, render_buffer, rng_state, + path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride); } } @@ -759,7 +756,6 @@ public: shader_kernel()(&kg, (uint4*)task.shader_input, (float4*)task.shader_output, - (float*)task.shader_output_luma, task.shader_eval_type, task.shader_filter, x, @@ -913,7 +909,6 @@ bool CPUSplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim, (void*)split_data.device_pointer, num_global_elements, (char*)ray_state.device_pointer, - (uint*)rtile.rng_state, rtile.start_sample, rtile.start_sample + rtile.num_samples, rtile.x, diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index a77212fb1fb..48ffa1484fb 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -1173,7 +1173,6 @@ public: task->rect.z-task->rect.x, task->rect.w-task->rect.y); - bool use_split_variance = use_split_kernel(); void *args[] = {&task->render_buffer.samples, &task->tiles_mem.device_pointer, &a_ptr, @@ -1183,8 +1182,7 @@ public: &buffer_variance_ptr, &task->rect, &task->render_buffer.pass_stride, - &task->render_buffer.denoising_data_offset, - &use_split_variance}; + &task->render_buffer.denoising_data_offset}; CUDA_LAUNCH_KERNEL(cuFilterDivideShadow, args); cuda_assert(cuCtxSynchronize()); @@ -1209,7 +1207,6 @@ public: task->rect.z-task->rect.x, task->rect.w-task->rect.y); - bool use_split_variance = use_split_kernel(); void *args[] = {&task->render_buffer.samples, &task->tiles_mem.device_pointer, &mean_offset, @@ -1218,8 +1215,7 @@ public: &variance_ptr, &task->rect, &task->render_buffer.pass_stride, - &task->render_buffer.denoising_data_offset, - &use_split_variance}; + &task->render_buffer.denoising_data_offset}; CUDA_LAUNCH_KERNEL(cuFilterGetFeature, args); cuda_assert(cuCtxSynchronize()); @@ -1285,19 +1281,16 @@ public: task.unmap_neighbor_tiles(rtiles, this); } - void path_trace(RenderTile& rtile, int sample, bool branched) + void path_trace(DeviceTask& task, RenderTile& rtile) { if(have_error()) return; CUDAContextScope scope(this); - CUfunction cuPathTrace; - CUdeviceptr d_buffer = cuda_device_ptr(rtile.buffer); - CUdeviceptr d_rng_state = cuda_device_ptr(rtile.rng_state); - /* get kernel function */ - if(branched) { + /* Get kernel function. */ + if(task.integrator_branched) { cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace")); } else { @@ -1308,40 +1301,65 @@ public: return; } - /* pass in parameters */ - void *args[] = {&d_buffer, - &d_rng_state, - &sample, - &rtile.x, - &rtile.y, - &rtile.w, - &rtile.h, - &rtile.offset, - &rtile.stride}; - - /* launch kernel */ - int threads_per_block; - cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuPathTrace)); - - /*int num_registers; - cuda_assert(cuFuncGetAttribute(&num_registers, CU_FUNC_ATTRIBUTE_NUM_REGS, cuPathTrace)); + cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1)); - printf("threads_per_block %d\n", threads_per_block); - printf("num_registers %d\n", num_registers);*/ + /* Allocate work tile. */ + device_vector<WorkTile> work_tiles; + work_tiles.resize(1); + + WorkTile *wtile = work_tiles.get_data(); + wtile->x = rtile.x; + wtile->y = rtile.y; + wtile->w = rtile.w; + wtile->h = rtile.h; + wtile->offset = rtile.offset; + wtile->stride = rtile.stride; + wtile->buffer = (float*)cuda_device_ptr(rtile.buffer); + mem_alloc("work_tiles", work_tiles, MEM_READ_ONLY); + + CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer); + + /* Prepare work size. More step samples render faster, but for now we + * remain conservative to avoid driver timeouts. */ + int min_blocks, num_threads_per_block; + cuda_assert(cuOccupancyMaxPotentialBlockSize(&min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0)); + uint step_samples = divide_up(min_blocks * num_threads_per_block, wtile->w * wtile->h);; + + /* Render all samples. */ + int start_sample = rtile.start_sample; + int end_sample = rtile.start_sample + rtile.num_samples; + + for(int sample = start_sample; sample < end_sample; sample += step_samples) { + /* Setup and copy work tile to device. */ + wtile->start_sample = sample; + wtile->num_samples = min(step_samples, end_sample - sample);; + mem_copy_to(work_tiles); + + uint total_work_size = wtile->w * wtile->h * wtile->num_samples; + uint num_blocks = divide_up(total_work_size, num_threads_per_block); + + /* Launch kernel. */ + void *args[] = {&d_work_tiles, + &total_work_size}; + + cuda_assert(cuLaunchKernel(cuPathTrace, + num_blocks, 1, 1, + num_threads_per_block, 1, 1, + 0, 0, args, 0)); - int xthreads = (int)sqrt(threads_per_block); - int ythreads = (int)sqrt(threads_per_block); - int xblocks = (rtile.w + xthreads - 1)/xthreads; - int yblocks = (rtile.h + ythreads - 1)/ythreads; + cuda_assert(cuCtxSynchronize()); - cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1)); + /* Update progress. */ + rtile.sample = sample + wtile->num_samples; + task.update_progress(&rtile, rtile.w*rtile.h*wtile->num_samples); - cuda_assert(cuLaunchKernel(cuPathTrace, - xblocks , yblocks, 1, /* blocks */ - xthreads, ythreads, 1, /* threads */ - 0, 0, args, 0)); + if(task.get_cancel()) { + if(task.need_finish_queue == false) + break; + } + } - cuda_assert(cuCtxSynchronize()); + mem_free(work_tiles); } void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half) @@ -1406,14 +1424,16 @@ public: CUfunction cuShader; CUdeviceptr d_input = cuda_device_ptr(task.shader_input); CUdeviceptr d_output = cuda_device_ptr(task.shader_output); - CUdeviceptr d_output_luma = cuda_device_ptr(task.shader_output_luma); /* get kernel function */ if(task.shader_eval_type >= SHADER_EVAL_BAKE) { cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_bake")); } + else if(task.shader_eval_type == SHADER_EVAL_DISPLACE) { + cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_displace")); + } else { - cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_shader")); + cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_background")); } /* do tasks in smaller chunks, so we can cancel it */ @@ -1432,9 +1452,6 @@ public: int arg = 0; args[arg++] = &d_input; args[arg++] = &d_output; - if(task.shader_eval_type < SHADER_EVAL_BAKE) { - args[arg++] = &d_output_luma; - } args[arg++] = &task.shader_eval_type; if(task.shader_eval_type >= SHADER_EVAL_BAKE) { args[arg++] = &task.shader_filter; @@ -1720,8 +1737,6 @@ public: if(task->type == DeviceTask::RENDER) { RenderTile tile; - bool branched = task->integrator_branched; - /* Upload Bindless Mapping */ load_bindless_mapping(); @@ -1745,21 +1760,7 @@ public: split_kernel->path_trace(task, tile, void_buffer, void_buffer); } else { - int start_sample = tile.start_sample; - int end_sample = tile.start_sample + tile.num_samples; - - for(int sample = start_sample; sample < end_sample; sample++) { - if(task->get_cancel()) { - if(task->need_finish_queue == false) - break; - } - - path_trace(tile, sample, branched); - - tile.sample = sample + 1; - - task->update_progress(&tile, tile.w*tile.h); - } + path_trace(*task, tile); } } else if(tile.task == RenderTile::DENOISE) { @@ -1960,7 +1961,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim CUdeviceptr d_use_queues_flag = device->cuda_device_ptr(use_queues_flag.device_pointer); CUdeviceptr d_work_pool_wgs = device->cuda_device_ptr(work_pool_wgs.device_pointer); - CUdeviceptr d_rng_state = device->cuda_device_ptr(rtile.rng_state); CUdeviceptr d_buffer = device->cuda_device_ptr(rtile.buffer); int end_sample = rtile.start_sample + rtile.num_samples; @@ -1970,7 +1970,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim CUdeviceptr* split_data_buffer; int* num_elements; CUdeviceptr* ray_state; - CUdeviceptr* rng_state; int* start_sample; int* end_sample; int* sx; @@ -1991,7 +1990,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim &d_split_data, &num_global_elements, &d_ray_state, - &d_rng_state, &rtile.start_sample, &end_sample, &rtile.x, diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h index b63dd00068b..20707ad04c9 100644 --- a/intern/cycles/device/device_memory.h +++ b/intern/cycles/device/device_memory.h @@ -46,6 +46,7 @@ enum MemoryType { /* Supported Data Types */ enum DataType { + TYPE_UNKNOWN, TYPE_UCHAR, TYPE_UINT, TYPE_INT, @@ -57,6 +58,7 @@ enum DataType { static inline size_t datatype_size(DataType datatype) { switch(datatype) { + case TYPE_UNKNOWN: return 1; case TYPE_UCHAR: return sizeof(uchar); case TYPE_FLOAT: return sizeof(float); case TYPE_UINT: return sizeof(uint); @@ -70,8 +72,8 @@ static inline size_t datatype_size(DataType datatype) /* Traits for data types */ template<typename T> struct device_type_traits { - static const DataType data_type = TYPE_UCHAR; - static const int num_elements = 0; + static const DataType data_type = TYPE_UNKNOWN; + static const int num_elements = sizeof(T); }; template<> struct device_type_traits<uchar> { diff --git a/intern/cycles/device/device_multi.cpp b/intern/cycles/device/device_multi.cpp index 35ae0303d6e..1b1a577d52f 100644 --- a/intern/cycles/device/device_multi.cpp +++ b/intern/cycles/device/device_multi.cpp @@ -284,7 +284,6 @@ public: foreach(SubDevice& sub, devices) { if(sub.device == sub_device) { if(tile.buffer) tile.buffer = sub.ptr_map[tile.buffer]; - if(tile.rng_state) tile.rng_state = sub.ptr_map[tile.rng_state]; } } } @@ -387,7 +386,6 @@ public: if(task.rgba_half) subtask.rgba_half = sub.ptr_map[task.rgba_half]; if(task.shader_input) subtask.shader_input = sub.ptr_map[task.shader_input]; if(task.shader_output) subtask.shader_output = sub.ptr_map[task.shader_output]; - if(task.shader_output_luma) subtask.shader_output_luma = sub.ptr_map[task.shader_output_luma]; sub.device->task_add(subtask); } diff --git a/intern/cycles/device/device_network.cpp b/intern/cycles/device/device_network.cpp index 571ba9465ca..deea59f1d23 100644 --- a/intern/cycles/device/device_network.cpp +++ b/intern/cycles/device/device_network.cpp @@ -660,10 +660,6 @@ protected: if(task.shader_output) task.shader_output = device_ptr_from_client_pointer(task.shader_output); - if(task.shader_output_luma) - task.shader_output_luma = device_ptr_from_client_pointer(task.shader_output_luma); - - task.acquire_tile = function_bind(&DeviceServer::task_acquire_tile, this, _1, _2); task.release_tile = function_bind(&DeviceServer::task_release_tile, this, _1); task.update_progress_sample = function_bind(&DeviceServer::task_update_progress_sample, this); @@ -737,7 +733,6 @@ protected: tile = entry.tile; if(tile.buffer) tile.buffer = ptr_map[tile.buffer]; - if(tile.rng_state) tile.rng_state = ptr_map[tile.rng_state]; result = true; break; @@ -769,7 +764,6 @@ protected: thread_scoped_lock acquire_lock(acquire_mutex); if(tile.buffer) tile.buffer = ptr_imap[tile.buffer]; - if(tile.rng_state) tile.rng_state = ptr_imap[tile.rng_state]; { thread_scoped_lock lock(rpc_lock); diff --git a/intern/cycles/device/device_network.h b/intern/cycles/device/device_network.h index a5d24c66018..3d3bd99dfe7 100644 --- a/intern/cycles/device/device_network.h +++ b/intern/cycles/device/device_network.h @@ -132,7 +132,7 @@ public: archive & type & task.x & task.y & task.w & task.h; archive & task.rgba_byte & task.rgba_half & task.buffer & task.sample & task.num_samples; archive & task.offset & task.stride; - archive & task.shader_input & task.shader_output & task.shader_output_luma & task.shader_eval_type; + archive & task.shader_input & task.shader_output & task.shader_eval_type; archive & task.shader_x & task.shader_w; archive & task.need_finish_queue; } @@ -142,7 +142,7 @@ public: archive & tile.x & tile.y & tile.w & tile.h; archive & tile.start_sample & tile.num_samples & tile.sample; archive & tile.resolution & tile.offset & tile.stride; - archive & tile.buffer & tile.rng_state; + archive & tile.buffer; } void write() @@ -291,7 +291,7 @@ public: *archive & type & task.x & task.y & task.w & task.h; *archive & task.rgba_byte & task.rgba_half & task.buffer & task.sample & task.num_samples; *archive & task.offset & task.stride; - *archive & task.shader_input & task.shader_output & task.shader_output_luma & task.shader_eval_type; + *archive & task.shader_input & task.shader_output & task.shader_eval_type; *archive & task.shader_x & task.shader_w; *archive & task.need_finish_queue; @@ -303,7 +303,7 @@ public: *archive & tile.x & tile.y & tile.w & tile.h; *archive & tile.start_sample & tile.num_samples & tile.sample; *archive & tile.resolution & tile.offset & tile.stride; - *archive & tile.buffer & tile.rng_state; + *archive & tile.buffer; tile.buffers = NULL; } diff --git a/intern/cycles/device/device_task.cpp b/intern/cycles/device/device_task.cpp index 3bc4c310283..3c7d24fb5b7 100644 --- a/intern/cycles/device/device_task.cpp +++ b/intern/cycles/device/device_task.cpp @@ -31,7 +31,7 @@ CCL_NAMESPACE_BEGIN DeviceTask::DeviceTask(Type type_) : type(type_), x(0), y(0), w(0), h(0), rgba_byte(0), rgba_half(0), buffer(0), sample(0), num_samples(1), - shader_input(0), shader_output(0), shader_output_luma(0), + shader_input(0), shader_output(0), shader_eval_type(0), shader_filter(0), shader_x(0), shader_w(0) { last_update_time = time_dt(); diff --git a/intern/cycles/device/device_task.h b/intern/cycles/device/device_task.h index 44a1efff1f5..b9658eb978f 100644 --- a/intern/cycles/device/device_task.h +++ b/intern/cycles/device/device_task.h @@ -46,7 +46,7 @@ public: int offset, stride; device_ptr shader_input; - device_ptr shader_output, shader_output_luma; + device_ptr shader_output; int shader_eval_type; int shader_filter; int shader_x, shader_w; diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp index 7bdf81462b8..3db3efd1103 100644 --- a/intern/cycles/device/opencl/opencl_base.cpp +++ b/intern/cycles/device/opencl/opencl_base.cpp @@ -228,7 +228,8 @@ bool OpenCLDeviceBase::load_kernels(const DeviceRequestedFeatures& requested_fea base_program = OpenCLProgram(this, "base", "kernel.cl", build_options_for_base_program(requested_features)); base_program.add_kernel(ustring("convert_to_byte")); base_program.add_kernel(ustring("convert_to_half_float")); - base_program.add_kernel(ustring("shader")); + base_program.add_kernel(ustring("displace")); + base_program.add_kernel(ustring("background")); base_program.add_kernel(ustring("bake")); base_program.add_kernel(ustring("zero_buffer")); @@ -982,7 +983,6 @@ bool OpenCLDeviceBase::denoising_divide_shadow(device_ptr a_ptr, cl_kernel ckFilterDivideShadow = denoising_program(ustring("filter_divide_shadow")); - char split_kernel = is_split_kernel()? 1 : 0; kernel_set_args(ckFilterDivideShadow, 0, task->render_buffer.samples, tiles_mem, @@ -993,8 +993,7 @@ bool OpenCLDeviceBase::denoising_divide_shadow(device_ptr a_ptr, buffer_variance_mem, task->rect, task->render_buffer.pass_stride, - task->render_buffer.denoising_data_offset, - split_kernel); + task->render_buffer.denoising_data_offset); enqueue_kernel(ckFilterDivideShadow, task->rect.z-task->rect.x, task->rect.w-task->rect.y); @@ -1015,7 +1014,6 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset, cl_kernel ckFilterGetFeature = denoising_program(ustring("filter_get_feature")); - char split_kernel = is_split_kernel()? 1 : 0; kernel_set_args(ckFilterGetFeature, 0, task->render_buffer.samples, tiles_mem, @@ -1025,8 +1023,7 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset, variance_mem, task->rect, task->render_buffer.pass_stride, - task->render_buffer.denoising_data_offset, - split_kernel); + task->render_buffer.denoising_data_offset); enqueue_kernel(ckFilterGetFeature, task->rect.z-task->rect.x, task->rect.w-task->rect.y); @@ -1116,7 +1113,6 @@ void OpenCLDeviceBase::shader(DeviceTask& task) cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer); cl_mem d_input = CL_MEM_PTR(task.shader_input); cl_mem d_output = CL_MEM_PTR(task.shader_output); - cl_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma); cl_int d_shader_eval_type = task.shader_eval_type; cl_int d_shader_filter = task.shader_filter; cl_int d_shader_x = task.shader_x; @@ -1125,10 +1121,15 @@ void OpenCLDeviceBase::shader(DeviceTask& task) cl_kernel kernel; - if(task.shader_eval_type >= SHADER_EVAL_BAKE) + if(task.shader_eval_type >= SHADER_EVAL_BAKE) { kernel = base_program(ustring("bake")); - else - kernel = base_program(ustring("shader")); + } + else if(task.shader_eval_type >= SHADER_EVAL_DISPLACE) { + kernel = base_program(ustring("displace")); + } + else { + kernel = base_program(ustring("background")); + } cl_uint start_arg_index = kernel_set_args(kernel, @@ -1137,12 +1138,6 @@ void OpenCLDeviceBase::shader(DeviceTask& task) d_input, d_output); - if(task.shader_eval_type < SHADER_EVAL_BAKE) { - start_arg_index += kernel_set_args(kernel, - start_arg_index, - d_output_luma); - } - set_kernel_arg_buffers(kernel, &start_arg_index); start_arg_index += kernel_set_args(kernel, diff --git a/intern/cycles/device/opencl/opencl_mega.cpp b/intern/cycles/device/opencl/opencl_mega.cpp index ec47fdafa3d..f4555eaba4f 100644 --- a/intern/cycles/device/opencl/opencl_mega.cpp +++ b/intern/cycles/device/opencl/opencl_mega.cpp @@ -62,7 +62,6 @@ public: /* Cast arguments to cl types. */ cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer); cl_mem d_buffer = CL_MEM_PTR(rtile.buffer); - cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state); cl_int d_x = rtile.x; cl_int d_y = rtile.y; cl_int d_w = rtile.w; @@ -79,8 +78,7 @@ public: kernel_set_args(ckPathTraceKernel, 0, d_data, - d_buffer, - d_rng_state); + d_buffer); set_kernel_arg_buffers(ckPathTraceKernel, &start_arg_index); diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index 16a96213100..976cc9df46d 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -192,7 +192,6 @@ 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; @@ -225,8 +224,7 @@ public: kg, data, *cached_memory.split_data, - *cached_memory.ray_state, - *cached_memory.rng_state); + *cached_memory.ray_state); device->set_kernel_arg_buffers(program(), &start_arg_index); @@ -356,8 +354,7 @@ public: kernel_data, split_data, num_global_elements, - ray_state, - rtile.rng_state); + ray_state); device->set_kernel_arg_buffers(device->program_data_init(), &start_arg_index); @@ -401,7 +398,6 @@ public: 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; diff --git a/intern/cycles/kernel/filter/filter_prefilter.h b/intern/cycles/kernel/filter/filter_prefilter.h index 2aeb54a62be..eefcbfea230 100644 --- a/intern/cycles/kernel/filter/filter_prefilter.h +++ b/intern/cycles/kernel/filter/filter_prefilter.h @@ -35,8 +35,7 @@ ccl_device void kernel_filter_divide_shadow(int sample, ccl_global float *bufferVariance, int4 rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int xtile = (x < tiles->x[1])? 0: ((x < tiles->x[2])? 1: 2); int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2); @@ -57,10 +56,12 @@ ccl_device void kernel_filter_divide_shadow(int sample, float varB = center_buffer[5]; int odd_sample = (sample+1)/2; int even_sample = sample/2; - if(use_split_variance) { - varA = max(0.0f, varA - unfilteredA[idx]*unfilteredA[idx]*odd_sample); - varB = max(0.0f, varB - unfilteredB[idx]*unfilteredB[idx]*even_sample); - } + + /* Approximate variance as E[x^2] - 1/N * (E[x])^2, since online variance + * update does not work efficiently with atomics in the kernel. */ + varA = max(0.0f, varA - unfilteredA[idx]*unfilteredA[idx]*odd_sample); + varB = max(0.0f, varB - unfilteredB[idx]*unfilteredB[idx]*even_sample); + varA /= max(odd_sample - 1, 1); varB /= max(even_sample - 1, 1); @@ -84,8 +85,7 @@ ccl_device void kernel_filter_get_feature(int sample, ccl_global float *mean, ccl_global float *variance, int4 rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int xtile = (x < tiles->x[1])? 0: ((x < tiles->x[2])? 1: 2); int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2); @@ -97,12 +97,9 @@ ccl_device void kernel_filter_get_feature(int sample, mean[idx] = center_buffer[m_offset] / sample; if(sample > 1) { - if(use_split_variance) { - variance[idx] = max(0.0f, (center_buffer[v_offset] - mean[idx]*mean[idx]*sample) / (sample * (sample-1))); - } - else { - variance[idx] = center_buffer[v_offset] / (sample * (sample-1)); - } + /* Approximate variance as E[x^2] - 1/N * (E[x])^2, since online variance + * update does not work efficiently with atomics in the kernel. */ + variance[idx] = max(0.0f, (center_buffer[v_offset] - mean[idx]*mean[idx]*sample) / (sample * (sample-1))); } else { /* Can't compute variance with single sample, just set it very high. */ diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h index 4d89839c46c..84d8d84d486 100644 --- a/intern/cycles/kernel/kernel_bake.h +++ b/intern/cycles/kernel/kernel_bake.h @@ -335,12 +335,14 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input, /* data passes */ case SHADER_EVAL_NORMAL: { + float3 N = sd.N; if((sd.flag & SD_HAS_BUMP)) { shader_eval_surface(kg, &sd, &state, 0); + N = shader_bsdf_average_normal(kg, &sd); } /* encoding: normal = (2 * color) - 1 */ - out = shader_bsdf_average_normal(kg, &sd) * 0.5f + make_float3(0.5f, 0.5f, 0.5f); + out = N * 0.5f + make_float3(0.5f, 0.5f, 0.5f); break; } case SHADER_EVAL_UV: @@ -491,78 +493,69 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input, #endif /* __BAKING__ */ -ccl_device void kernel_shader_evaluate(KernelGlobals *kg, - ccl_global uint4 *input, - ccl_global float4 *output, - ccl_global float *output_luma, - ShaderEvalType type, - int i, - int sample) +ccl_device void kernel_displace_evaluate(KernelGlobals *kg, + ccl_global uint4 *input, + ccl_global float4 *output, + int i) { ShaderData sd; PathState state = {0}; uint4 in = input[i]; - float3 out; - if(type == SHADER_EVAL_DISPLACE) { - /* setup shader data */ - int object = in.x; - int prim = in.y; - float u = __uint_as_float(in.z); - float v = __uint_as_float(in.w); + /* setup shader data */ + int object = in.x; + int prim = in.y; + float u = __uint_as_float(in.z); + float v = __uint_as_float(in.w); - shader_setup_from_displace(kg, &sd, object, prim, u, v); + shader_setup_from_displace(kg, &sd, object, prim, u, v); - /* evaluate */ - float3 P = sd.P; - shader_eval_displacement(kg, &sd, &state); - out = sd.P - P; + /* evaluate */ + float3 P = sd.P; + shader_eval_displacement(kg, &sd, &state); + float3 D = sd.P - P; - object_inverse_dir_transform(kg, &sd, &out); - } - else { // SHADER_EVAL_BACKGROUND - /* setup ray */ - Ray ray; - float u = __uint_as_float(in.x); - float v = __uint_as_float(in.y); - - ray.P = make_float3(0.0f, 0.0f, 0.0f); - ray.D = equirectangular_to_direction(u, v); - ray.t = 0.0f; + object_inverse_dir_transform(kg, &sd, &D); + + /* write output */ + output[i] += make_float4(D.x, D.y, D.z, 0.0f); +} + +ccl_device void kernel_background_evaluate(KernelGlobals *kg, + ccl_global uint4 *input, + ccl_global float4 *output, + int i) +{ + ShaderData sd; + PathState state = {0}; + uint4 in = input[i]; + + /* setup ray */ + Ray ray; + float u = __uint_as_float(in.x); + float v = __uint_as_float(in.y); + + ray.P = make_float3(0.0f, 0.0f, 0.0f); + ray.D = equirectangular_to_direction(u, v); + ray.t = 0.0f; #ifdef __CAMERA_MOTION__ - ray.time = 0.5f; + ray.time = 0.5f; #endif #ifdef __RAY_DIFFERENTIALS__ - ray.dD = differential3_zero(); - ray.dP = differential3_zero(); + ray.dD = differential3_zero(); + ray.dP = differential3_zero(); #endif - /* setup shader data */ - shader_setup_from_background(kg, &sd, &ray); + /* setup shader data */ + shader_setup_from_background(kg, &sd, &ray); + + /* evaluate */ + int flag = 0; /* we can't know which type of BSDF this is for */ + float3 color = shader_eval_background(kg, &sd, &state, flag); - /* evaluate */ - int flag = 0; /* we can't know which type of BSDF this is for */ - out = shader_eval_background(kg, &sd, &state, flag); - } - /* write output */ - if(sample == 0) { - if(output != NULL) { - output[i] = make_float4(out.x, out.y, out.z, 0.0f); - } - if(output_luma != NULL) { - output_luma[i] = average(out); - } - } - else { - if(output != NULL) { - output[i] += make_float4(out.x, out.y, out.z, 0.0f); - } - if(output_luma != NULL) { - output_luma[i] += average(out); - } - } + output[i] += make_float4(color.x, color.y, color.z, 0.0f); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index 1e2af9de8b3..2e8ca48c413 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -50,7 +50,8 @@ #endif #define ccl_device_noinline __device__ __noinline__ #define ccl_global -#define ccl_constant +#define ccl_static_constant __constant__ +#define ccl_constant const #define ccl_local __shared__ #define ccl_local_param #define ccl_private diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h index 36d6031d042..7f81523791b 100644 --- a/intern/cycles/kernel/kernel_compat_opencl.h +++ b/intern/cycles/kernel/kernel_compat_opencl.h @@ -36,6 +36,7 @@ #define ccl_device_forceinline ccl_device #define ccl_device_noinline ccl_device ccl_noinline #define ccl_may_alias +#define ccl_static_constant static __constant #define ccl_constant __constant #define ccl_global __global #define ccl_local __local diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h index fff7f4cfdb7..644cc173571 100644 --- a/intern/cycles/kernel/kernel_passes.h +++ b/intern/cycles/kernel/kernel_passes.h @@ -16,19 +16,23 @@ CCL_NAMESPACE_BEGIN -ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, int sample, float value) +#if defined(__SPLIT_KERNEL__) || defined(__KERNEL_CUDA__) +#define __ATOMIC_PASS_WRITE__ +#endif + +ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, float value) { ccl_global float *buf = buffer; -#if defined(__SPLIT_KERNEL__) +#ifdef __ATOMIC_PASS_WRITE__ atomic_add_and_fetch_float(buf, value); #else - *buf = (sample == 0)? value: *buf + value; -#endif /* __SPLIT_KERNEL__ */ + *buf += value; +#endif } -ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, int sample, float3 value) +ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, float3 value) { -#if defined(__SPLIT_KERNEL__) +#ifdef __ATOMIC_PASS_WRITE__ ccl_global float *buf_x = buffer + 0; ccl_global float *buf_y = buffer + 1; ccl_global float *buf_z = buffer + 2; @@ -38,13 +42,13 @@ ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, int sa atomic_add_and_fetch_float(buf_z, value.z); #else ccl_global float3 *buf = (ccl_global float3*)buffer; - *buf = (sample == 0)? value: *buf + value; -#endif /* __SPLIT_KERNEL__ */ + *buf += value; +#endif } -ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, int sample, float4 value) +ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, float4 value) { -#if defined(__SPLIT_KERNEL__) +#ifdef __ATOMIC_PASS_WRITE__ ccl_global float *buf_x = buffer + 0; ccl_global float *buf_y = buffer + 1; ccl_global float *buf_z = buffer + 2; @@ -56,58 +60,35 @@ ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, int sa atomic_add_and_fetch_float(buf_w, value.w); #else ccl_global float4 *buf = (ccl_global float4*)buffer; - *buf = (sample == 0)? value: *buf + value; -#endif /* __SPLIT_KERNEL__ */ + *buf += value; +#endif } #ifdef __DENOISING_FEATURES__ -ccl_device_inline void kernel_write_pass_float_variance(ccl_global float *buffer, int sample, float value) +ccl_device_inline void kernel_write_pass_float_variance(ccl_global float *buffer, float value) { - kernel_write_pass_float(buffer, sample, value); + kernel_write_pass_float(buffer, value); /* The online one-pass variance update that's used for the megakernel can't easily be implemented * with atomics, so for the split kernel the E[x^2] - 1/N * (E[x])^2 fallback is used. */ -# ifdef __SPLIT_KERNEL__ - kernel_write_pass_float(buffer+1, sample, value*value); -# else - if(sample == 0) { - kernel_write_pass_float(buffer+1, sample, 0.0f); - } - else { - float new_mean = buffer[0] * (1.0f / (sample + 1)); - float old_mean = (buffer[0] - value) * (1.0f / sample); - kernel_write_pass_float(buffer+1, sample, (value - new_mean) * (value - old_mean)); - } -# endif + kernel_write_pass_float(buffer+1, value*value); } -# if defined(__SPLIT_KERNEL__) +# ifdef __ATOMIC_PASS_WRITE__ # define kernel_write_pass_float3_unaligned kernel_write_pass_float3 # else -ccl_device_inline void kernel_write_pass_float3_unaligned(ccl_global float *buffer, int sample, float3 value) +ccl_device_inline void kernel_write_pass_float3_unaligned(ccl_global float *buffer, float3 value) { - buffer[0] = (sample == 0)? value.x: buffer[0] + value.x; - buffer[1] = (sample == 0)? value.y: buffer[1] + value.y; - buffer[2] = (sample == 0)? value.z: buffer[2] + value.z; + buffer[0] += value.x; + buffer[1] += value.y; + buffer[2] += value.z; } # endif -ccl_device_inline void kernel_write_pass_float3_variance(ccl_global float *buffer, int sample, float3 value) +ccl_device_inline void kernel_write_pass_float3_variance(ccl_global float *buffer, float3 value) { - kernel_write_pass_float3_unaligned(buffer, sample, value); -# ifdef __SPLIT_KERNEL__ - kernel_write_pass_float3_unaligned(buffer+3, sample, value*value); -# else - if(sample == 0) { - kernel_write_pass_float3_unaligned(buffer+3, sample, make_float3(0.0f, 0.0f, 0.0f)); - } - else { - float3 sum = make_float3(buffer[0], buffer[1], buffer[2]); - float3 new_mean = sum * (1.0f / (sample + 1)); - float3 old_mean = (sum - value) * (1.0f / sample); - kernel_write_pass_float3_unaligned(buffer+3, sample, (value - new_mean) * (value - old_mean)); - } -# endif + kernel_write_pass_float3_unaligned(buffer, value); + kernel_write_pass_float3_unaligned(buffer+3, value*value); } ccl_device_inline void kernel_write_denoising_shadow(KernelGlobals *kg, ccl_global float *buffer, @@ -121,22 +102,11 @@ ccl_device_inline void kernel_write_denoising_shadow(KernelGlobals *kg, ccl_glob path_total = ensure_finite(path_total); path_total_shaded = ensure_finite(path_total_shaded); - kernel_write_pass_float(buffer, sample/2, path_total); - kernel_write_pass_float(buffer+1, sample/2, path_total_shaded); + kernel_write_pass_float(buffer, path_total); + kernel_write_pass_float(buffer+1, path_total_shaded); float value = path_total_shaded / max(path_total, 1e-7f); -# ifdef __SPLIT_KERNEL__ - kernel_write_pass_float(buffer+2, sample/2, value*value); -# else - if(sample < 2) { - kernel_write_pass_float(buffer+2, sample/2, 0.0f); - } - else { - float old_value = (buffer[1] - path_total_shaded) / max(buffer[0] - path_total, 1e-7f); - float new_value = buffer[1] / max(buffer[0], 1e-7f); - kernel_write_pass_float(buffer+2, sample, (value - new_value) * (value - old_value)); - } -# endif + kernel_write_pass_float(buffer+2, value*value); } #endif /* __DENOISING_FEATURES__ */ @@ -197,28 +167,23 @@ ccl_device_inline void kernel_update_denoising_features(KernelGlobals *kg, #ifdef __KERNEL_DEBUG__ ccl_device_inline void kernel_write_debug_passes(KernelGlobals *kg, ccl_global float *buffer, - PathRadiance *L, - int sample) + PathRadiance *L) { int flag = kernel_data.film.pass_flag; if(flag & PASS_BVH_TRAVERSED_NODES) { kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_nodes, - sample, L->debug_data.num_bvh_traversed_nodes); } if(flag & PASS_BVH_TRAVERSED_INSTANCES) { kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_instances, - sample, L->debug_data.num_bvh_traversed_instances); } if(flag & PASS_BVH_INTERSECTIONS) { kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_intersections, - sample, L->debug_data.num_bvh_intersections); } if(flag & PASS_RAY_BOUNCES) { kernel_write_pass_float(buffer + kernel_data.film.pass_ray_bounces, - sample, L->debug_data.num_ray_bounces); } } @@ -243,35 +208,33 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global fl kernel_data.film.pass_alpha_threshold == 0.0f || average(shader_bsdf_alpha(kg, sd)) >= kernel_data.film.pass_alpha_threshold) { - int sample = state->sample; - - if(sample == 0) { + if(state->sample == 0) { if(flag & PASS_DEPTH) { float depth = camera_distance(kg, sd->P); - kernel_write_pass_float(buffer + kernel_data.film.pass_depth, sample, depth); + kernel_write_pass_float(buffer + kernel_data.film.pass_depth, depth); } if(flag & PASS_OBJECT_ID) { float id = object_pass_id(kg, sd->object); - kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, sample, id); + kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, id); } if(flag & PASS_MATERIAL_ID) { float id = shader_pass_id(kg, sd); - kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, sample, id); + kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, id); } } if(flag & PASS_NORMAL) { float3 normal = shader_bsdf_average_normal(kg, sd); - kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, sample, normal); + kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, normal); } if(flag & PASS_UV) { float3 uv = primitive_uv(kg, sd); - kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, sample, uv); + kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, uv); } if(flag & PASS_MOTION) { float4 speed = primitive_motion_vector(kg, sd); - kernel_write_pass_float4(buffer + kernel_data.film.pass_motion, sample, speed); - kernel_write_pass_float(buffer + kernel_data.film.pass_motion_weight, sample, 1.0f); + kernel_write_pass_float4(buffer + kernel_data.film.pass_motion, speed); + kernel_write_pass_float(buffer + kernel_data.film.pass_motion_weight, 1.0f); } state->flag |= PATH_RAY_SINGLE_PASS_DONE; @@ -314,7 +277,7 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global fl #endif } -ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L, int sample) +ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L) { #ifdef __PASSES__ int flag = kernel_data.film.pass_flag; @@ -323,116 +286,90 @@ ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global f return; if(flag & PASS_DIFFUSE_INDIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, sample, L->indirect_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, L->indirect_diffuse); if(flag & PASS_GLOSSY_INDIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, sample, L->indirect_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, L->indirect_glossy); if(flag & PASS_TRANSMISSION_INDIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, sample, L->indirect_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, L->indirect_transmission); if(flag & PASS_SUBSURFACE_INDIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_indirect, sample, L->indirect_subsurface); + kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_indirect, L->indirect_subsurface); if(flag & PASS_DIFFUSE_DIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, sample, L->direct_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, L->direct_diffuse); if(flag & PASS_GLOSSY_DIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, sample, L->direct_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, L->direct_glossy); if(flag & PASS_TRANSMISSION_DIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, sample, L->direct_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, L->direct_transmission); if(flag & PASS_SUBSURFACE_DIRECT) - kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_direct, sample, L->direct_subsurface); + kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_direct, L->direct_subsurface); if(flag & PASS_EMISSION) - kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, sample, L->emission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, L->emission); if(flag & PASS_BACKGROUND) - kernel_write_pass_float3(buffer + kernel_data.film.pass_background, sample, L->background); + kernel_write_pass_float3(buffer + kernel_data.film.pass_background, L->background); if(flag & PASS_AO) - kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, sample, L->ao); + kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, L->ao); if(flag & PASS_DIFFUSE_COLOR) - kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, sample, L->color_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, L->color_diffuse); if(flag & PASS_GLOSSY_COLOR) - kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, sample, L->color_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, L->color_glossy); if(flag & PASS_TRANSMISSION_COLOR) - kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, sample, L->color_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, L->color_transmission); if(flag & PASS_SUBSURFACE_COLOR) - kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_color, sample, L->color_subsurface); + kernel_write_pass_float3(buffer + kernel_data.film.pass_subsurface_color, L->color_subsurface); if(flag & PASS_SHADOW) { float4 shadow = L->shadow; shadow.w = kernel_data.film.pass_shadow_scale; - kernel_write_pass_float4(buffer + kernel_data.film.pass_shadow, sample, shadow); + kernel_write_pass_float4(buffer + kernel_data.film.pass_shadow, shadow); } if(flag & PASS_MIST) - kernel_write_pass_float(buffer + kernel_data.film.pass_mist, sample, 1.0f - L->mist); + kernel_write_pass_float(buffer + kernel_data.film.pass_mist, 1.0f - L->mist); #endif } ccl_device_inline void kernel_write_result(KernelGlobals *kg, ccl_global float *buffer, int sample, PathRadiance *L) { - if(L) { - float alpha; - float3 L_sum = path_radiance_clamp_and_sum(kg, L, &alpha); + float alpha; + float3 L_sum = path_radiance_clamp_and_sum(kg, L, &alpha); - kernel_write_pass_float4(buffer, sample, make_float4(L_sum.x, L_sum.y, L_sum.z, alpha)); + kernel_write_pass_float4(buffer, make_float4(L_sum.x, L_sum.y, L_sum.z, alpha)); - kernel_write_light_passes(kg, buffer, L, sample); + kernel_write_light_passes(kg, buffer, L); #ifdef __DENOISING_FEATURES__ - if(kernel_data.film.pass_denoising_data) { + if(kernel_data.film.pass_denoising_data) { # ifdef __SHADOW_TRICKS__ - kernel_write_denoising_shadow(kg, buffer + kernel_data.film.pass_denoising_data, sample, average(L->path_total), average(L->path_total_shaded)); + kernel_write_denoising_shadow(kg, buffer + kernel_data.film.pass_denoising_data, sample, average(L->path_total), average(L->path_total_shaded)); # else - kernel_write_denoising_shadow(kg, buffer + kernel_data.film.pass_denoising_data, sample, 0.0f, 0.0f); + kernel_write_denoising_shadow(kg, buffer + kernel_data.film.pass_denoising_data, sample, 0.0f, 0.0f); # endif - if(kernel_data.film.pass_denoising_clean) { - float3 noisy, clean; - path_radiance_split_denoising(kg, L, &noisy, &clean); - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, - sample, noisy); - kernel_write_pass_float3_unaligned(buffer + kernel_data.film.pass_denoising_clean, - sample, clean); - } - else { - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, - sample, ensure_finite3(L_sum)); - } - - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_NORMAL, - sample, L->denoising_normal); - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_ALBEDO, - sample, L->denoising_albedo); - kernel_write_pass_float_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH, - sample, L->denoising_depth); + if(kernel_data.film.pass_denoising_clean) { + float3 noisy, clean; + path_radiance_split_denoising(kg, L, &noisy, &clean); + kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, + noisy); + kernel_write_pass_float3_unaligned(buffer + kernel_data.film.pass_denoising_clean, + clean); } + else { + kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, + ensure_finite3(L_sum)); + } + + kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_NORMAL, + L->denoising_normal); + kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_ALBEDO, + L->denoising_albedo); + kernel_write_pass_float_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH, + L->denoising_depth); + } #endif /* __DENOISING_FEATURES__ */ #ifdef __KERNEL_DEBUG__ - kernel_write_debug_passes(kg, buffer, L, sample); + kernel_write_debug_passes(kg, buffer, L); #endif - } - else { - kernel_write_pass_float4(buffer, sample, make_float4(0.0f, 0.0f, 0.0f, 0.0f)); - -#ifdef __DENOISING_FEATURES__ - if(kernel_data.film.pass_denoising_data) { - kernel_write_denoising_shadow(kg, buffer + kernel_data.film.pass_denoising_data, sample, 0.0f, 0.0f); - - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, - sample, make_float3(0.0f, 0.0f, 0.0f)); - - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_NORMAL, - sample, make_float3(0.0f, 0.0f, 0.0f)); - kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_ALBEDO, - sample, make_float3(0.0f, 0.0f, 0.0f)); - kernel_write_pass_float_variance(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH, - sample, 0.0f); - - if(kernel_data.film.pass_denoising_clean) { - kernel_write_pass_float3_unaligned(buffer + kernel_data.film.pass_denoising_clean, - sample, make_float3(0.0f, 0.0f, 0.0f)); - } - } -#endif /* __DENOISING_FEATURES__ */ - } } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 793fede0deb..6b6c5603b70 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -672,24 +672,22 @@ ccl_device_forceinline void kernel_path_integrate( } ccl_device void kernel_path_trace(KernelGlobals *kg, - ccl_global float *buffer, ccl_global uint *rng_state, + ccl_global float *buffer, int sample, int x, int y, int offset, int stride) { /* buffer offset */ int index = offset + x + y*stride; int pass_stride = kernel_data.film.pass_stride; - rng_state += index; buffer += index*pass_stride; /* Initialize random numbers and sample ray. */ uint rng_hash; Ray ray; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, &ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray); if(ray.t == 0.0f) { - kernel_write_result(kg, buffer, sample, NULL); return; } diff --git a/intern/cycles/kernel/kernel_path_branched.h b/intern/cycles/kernel/kernel_path_branched.h index 6e0ec22d581..2597d684a36 100644 --- a/intern/cycles/kernel/kernel_path_branched.h +++ b/intern/cycles/kernel/kernel_path_branched.h @@ -538,21 +538,20 @@ ccl_device void kernel_branched_path_integrate(KernelGlobals *kg, } ccl_device void kernel_branched_path_trace(KernelGlobals *kg, - ccl_global float *buffer, ccl_global uint *rng_state, + ccl_global float *buffer, int sample, int x, int y, int offset, int stride) { /* buffer offset */ int index = offset + x + y*stride; int pass_stride = kernel_data.film.pass_stride; - rng_state += index; buffer += index*pass_stride; /* initialize random numbers and ray */ uint rng_hash; Ray ray; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, &ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray); /* integrate */ PathRadiance L; @@ -561,9 +560,6 @@ ccl_device void kernel_branched_path_trace(KernelGlobals *kg, kernel_branched_path_integrate(kg, rng_hash, sample, ray, buffer, &L); kernel_write_result(kg, buffer, sample, &L); } - else { - kernel_write_result(kg, buffer, sample, NULL); - } } #endif /* __SPLIT_KERNEL__ */ diff --git a/intern/cycles/kernel/kernel_path_common.h b/intern/cycles/kernel/kernel_path_common.h index 54dd278a185..d83fd474cde 100644 --- a/intern/cycles/kernel/kernel_path_common.h +++ b/intern/cycles/kernel/kernel_path_common.h @@ -19,7 +19,6 @@ CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, - ccl_global uint *rng_state, int sample, int x, int y, uint *rng_hash, @@ -30,11 +29,7 @@ ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, int num_samples = kernel_data.integrator.aa_samples; - if(sample == kernel_data.integrator.start_sample) { - *rng_state = hash_int_2d(x, y); - } - - path_rng_init(kg, rng_state, sample, num_samples, rng_hash, x, y, &filter_u, &filter_v); + path_rng_init(kg, sample, num_samples, rng_hash, x, y, &filter_u, &filter_v); /* sample camera ray */ diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h index 11798d87cb5..e7a6134b8eb 100644 --- a/intern/cycles/kernel/kernel_random.h +++ b/intern/cycles/kernel/kernel_random.h @@ -15,6 +15,7 @@ */ #include "kernel/kernel_jitter.h" +#include "util/util_hash.h" CCL_NAMESPACE_BEGIN @@ -115,14 +116,13 @@ ccl_device_forceinline void path_rng_2D(KernelGlobals *kg, } ccl_device_inline void path_rng_init(KernelGlobals *kg, - ccl_global uint *rng_state, int sample, int num_samples, uint *rng_hash, int x, int y, float *fx, float *fy) { /* load state */ - *rng_hash = *rng_state; + *rng_hash = hash_int_2d(x, y); *rng_hash ^= kernel_data.integrator.seed; #ifdef __DEBUG_CORRELATION__ diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index eeb4eb0097f..695d4fc380a 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -1204,7 +1204,7 @@ ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ccl_ #ifdef __SVM__ # ifdef __OSL__ if(kg->osl) - OSLShader::eval_displacement(kg, sd); + OSLShader::eval_displacement(kg, sd, state); else # endif { diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index 6c5b6ca3b2d..19c77c1ed4f 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -1448,6 +1448,20 @@ enum RayState { #define PATCH_MAP_NODE_IS_LEAF (1u << 31) #define PATCH_MAP_NODE_INDEX_MASK (~(PATCH_MAP_NODE_IS_SET | PATCH_MAP_NODE_IS_LEAF)) +/* Work Tiles */ + +typedef struct WorkTile { + uint x, y, w, h; + + uint start_sample; + uint num_samples; + + uint offset; + uint stride; + + ccl_global float *buffer; +} WorkTile; + CCL_NAMESPACE_END #endif /* __KERNEL_TYPES_H__ */ diff --git a/intern/cycles/kernel/kernel_work_stealing.h b/intern/cycles/kernel/kernel_work_stealing.h index 0c11158e8da..0c2d9379b63 100644 --- a/intern/cycles/kernel/kernel_work_stealing.h +++ b/intern/cycles/kernel/kernel_work_stealing.h @@ -27,29 +27,28 @@ CCL_NAMESPACE_BEGIN # pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable #endif +#ifdef __SPLIT_KERNEL__ /* Returns true if there is work */ ccl_device bool get_next_work(KernelGlobals *kg, - uint thread_index, + ccl_global uint *work_pools, + uint total_work_size, + uint ray_index, ccl_private uint *global_work_index) { - uint total_work_size = kernel_split_params.w - * kernel_split_params.h - * kernel_split_params.num_samples; - /* With a small amount of work there may be more threads than work due to * rounding up of global size, stop such threads immediately. */ - if(thread_index >= total_work_size) { + if(ray_index >= total_work_size) { return false; } /* Increase atomic work index counter in pool. */ - uint pool = thread_index / WORK_POOL_SIZE; - uint work_index = atomic_fetch_and_inc_uint32(&kernel_split_params.work_pools[pool]); + uint pool = ray_index / WORK_POOL_SIZE; + uint work_index = atomic_fetch_and_inc_uint32(&work_pools[pool]); /* Map per-pool work index to a global work index. */ uint global_size = ccl_global_size(0) * ccl_global_size(1); kernel_assert(global_size % WORK_POOL_SIZE == 0); - kernel_assert(thread_index < global_size); + kernel_assert(ray_index < global_size); *global_work_index = (work_index / WORK_POOL_SIZE) * global_size + (pool * WORK_POOL_SIZE) @@ -58,23 +57,24 @@ ccl_device bool get_next_work(KernelGlobals *kg, /* Test if all work for this pool is done. */ return (*global_work_index < total_work_size); } +#endif -/* Map global work index to pixel X/Y and sample. */ -ccl_device_inline void get_work_pixel(KernelGlobals *kg, +/* Map global work index to tile, pixel X/Y and sample. */ +ccl_device_inline void get_work_pixel(ccl_global const WorkTile *tile, uint global_work_index, ccl_private uint *x, ccl_private uint *y, ccl_private uint *sample) { - uint tile_pixels = kernel_split_params.w * kernel_split_params.h; + uint tile_pixels = tile->w * tile->h; uint sample_offset = global_work_index / tile_pixels; uint pixel_offset = global_work_index - sample_offset * tile_pixels; - uint y_offset = pixel_offset / kernel_split_params.w; - uint x_offset = pixel_offset - y_offset * kernel_split_params.w; + uint y_offset = pixel_offset / tile->w; + uint x_offset = pixel_offset - y_offset * tile->w; - *x = kernel_split_params.x + x_offset; - *y = kernel_split_params.y + y_offset; - *sample = kernel_split_params.start_sample + sample_offset; + *x = tile->x + x_offset; + *y = tile->y + y_offset; + *sample = tile->start_sample + sample_offset; } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu.h b/intern/cycles/kernel/kernels/cpu/filter_cpu.h index 2ed713299fd..bf13ba62806 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu.h @@ -27,8 +27,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, float *bufferV, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance); + int buffer_denoising_offset); void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, TilesInfo *tiles, @@ -40,8 +39,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, float *variance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance); + int buffer_denoising_offset); void KERNEL_FUNCTION_FULL_NAME(filter_detect_outliers)(int x, int y, ccl_global float *image, diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h index 8dc1a8d583c..2fbb0ea2bdb 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h @@ -45,8 +45,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, float *bufferVariance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_divide_shadow); @@ -60,8 +59,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, bufferVariance, load_int4(prefilter_rect), buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); #endif } @@ -74,8 +72,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, float *mean, float *variance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_get_feature); @@ -86,8 +83,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, mean, variance, load_int4(prefilter_rect), buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); #endif } diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h index c8938534fe8..6bdb8546a24 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h @@ -18,7 +18,6 @@ void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg, float *buffer, - unsigned int *rng_state, int sample, int x, int y, int offset, @@ -42,7 +41,6 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg, void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, uint4 *input, float4 *output, - float *output_luma, int type, int filter, int i, @@ -57,7 +55,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( ccl_global void *split_data_buffer, int num_elements, ccl_global char *ray_state, - ccl_global uint *rng_state, int start_sample, int end_sample, int sx, int sy, int sw, int sh, int offset, int stride, diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h index d4315ee5ec4..fdeb7dcd3e4 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h @@ -75,7 +75,6 @@ CCL_NAMESPACE_BEGIN void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg, float *buffer, - unsigned int *rng_state, int sample, int x, int y, int offset, @@ -88,7 +87,6 @@ void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg, if(kernel_data.integrator.branched) { kernel_branched_path_trace(kg, buffer, - rng_state, sample, x, y, offset, @@ -97,7 +95,7 @@ void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg, else # endif { - kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride); + kernel_path_trace(kg, buffer, sample, x, y, offset, stride); } #endif /* KERNEL_STUB */ } @@ -151,7 +149,6 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg, void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, uint4 *input, float4 *output, - float *output_luma, int type, int filter, int i, @@ -162,7 +159,6 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, STUB_ASSERT(KERNEL_ARCH, shader); #else if(type >= SHADER_EVAL_BAKE) { - kernel_assert(output_luma == NULL); # ifdef __BAKING__ kernel_bake_evaluate(kg, input, @@ -174,14 +170,11 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, sample); # endif } + else if(type == SHADER_EVAL_DISPLACE) { + kernel_displace_evaluate(kg, input, output, i); + } else { - kernel_shader_evaluate(kg, - input, - output, - output_luma, - (ShaderEvalType)type, - i, - sample); + kernel_background_evaluate(kg, input, output, i); } #endif /* KERNEL_STUB */ } diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu index 009c3fde9d5..c8172355a7f 100644 --- a/intern/cycles/kernel/kernels/cuda/filter.cu +++ b/intern/cycles/kernel/kernels/cuda/filter.cu @@ -37,8 +37,7 @@ kernel_cuda_filter_divide_shadow(int sample, float *bufferVariance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x; int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y; @@ -53,8 +52,7 @@ kernel_cuda_filter_divide_shadow(int sample, bufferVariance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } @@ -68,8 +66,7 @@ kernel_cuda_filter_get_feature(int sample, float *variance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x; int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y; @@ -81,8 +78,7 @@ kernel_cuda_filter_get_feature(int sample, mean, variance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu index dc343cb387a..1ac6afd167a 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel.cu @@ -20,6 +20,9 @@ #include "kernel/kernel_compat_cuda.h" #include "kernel_config.h" + +#include "util/util_atomic.h" + #include "kernel/kernel_math.h" #include "kernel/kernel_types.h" #include "kernel/kernel_globals.h" @@ -27,32 +30,37 @@ #include "kernel/kernel_path.h" #include "kernel/kernel_path_branched.h" #include "kernel/kernel_bake.h" +#include "kernel/kernel_work_stealing.h" /* kernels */ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride) +kernel_cuda_path_trace(WorkTile *tile, uint total_work_size) { - int x = sx + blockDim.x*blockIdx.x + threadIdx.x; - int y = sy + blockDim.y*blockIdx.y + threadIdx.y; + int work_index = ccl_global_id(0); + + if(work_index < total_work_size) { + uint x, y, sample; + get_work_pixel(tile, work_index, &x, &y, &sample); - if(x < sx + sw && y < sy + sh) { KernelGlobals kg; - kernel_path_trace(&kg, buffer, rng_state, sample, x, y, offset, stride); + kernel_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride); } } #ifdef __BRANCHED_PATH__ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_BRANCHED_MAX_REGISTERS) -kernel_cuda_branched_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride) +kernel_cuda_branched_path_trace(WorkTile *tile, uint total_work_size) { - int x = sx + blockDim.x*blockIdx.x + threadIdx.x; - int y = sy + blockDim.y*blockIdx.y + threadIdx.y; + int work_index = ccl_global_id(0); + + if(work_index < total_work_size) { + uint x, y, sample; + get_work_pixel(tile, work_index, &x, &y, &sample); - if(x < sx + sw && y < sy + sh) { KernelGlobals kg; - kernel_branched_path_trace(&kg, buffer, rng_state, sample, x, y, offset, stride); + kernel_branched_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride); } } #endif @@ -83,26 +91,37 @@ kernel_cuda_convert_to_half_float(uchar4 *rgba, float *buffer, float sample_scal extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_shader(uint4 *input, - float4 *output, - float *output_luma, - int type, - int sx, - int sw, - int offset, - int sample) +kernel_cuda_displace(uint4 *input, + float4 *output, + int type, + int sx, + int sw, + int offset, + int sample) +{ + int x = sx + blockDim.x*blockIdx.x + threadIdx.x; + + if(x < sx + sw) { + KernelGlobals kg; + kernel_displace_evaluate(&kg, input, output, x); + } +} + +extern "C" __global__ void +CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) +kernel_cuda_background(uint4 *input, + float4 *output, + int type, + int sx, + int sw, + int offset, + int sample) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; if(x < sx + sw) { KernelGlobals kg; - kernel_shader_evaluate(&kg, - input, - output, - output_luma, - (ShaderEvalType)type, - x, - sample); + kernel_background_evaluate(&kg, input, output, x); } } diff --git a/intern/cycles/kernel/kernels/cuda/kernel_split.cu b/intern/cycles/kernel/kernels/cuda/kernel_split.cu index e97e87285a5..43b3d0aa0e6 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel_split.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel_split.cu @@ -60,7 +60,6 @@ kernel_cuda_path_trace_data_init( ccl_global void *split_data_buffer, int num_elements, ccl_global char *ray_state, - ccl_global uint *rng_state, int start_sample, int end_sample, int sx, int sy, int sw, int sh, int offset, int stride, @@ -76,7 +75,6 @@ kernel_cuda_path_trace_data_init( split_data_buffer, num_elements, ray_state, - rng_state, start_sample, end_sample, sx, sy, sw, sh, offset, stride, diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl index f015ac47d8a..7a7b596a350 100644 --- a/intern/cycles/kernel/kernels/opencl/filter.cl +++ b/intern/cycles/kernel/kernels/opencl/filter.cl @@ -31,8 +31,7 @@ __kernel void kernel_ocl_filter_divide_shadow(int sample, ccl_global float *bufferVariance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - char use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + get_global_id(0); int y = prefilter_rect.y + get_global_id(1); @@ -47,8 +46,7 @@ __kernel void kernel_ocl_filter_divide_shadow(int sample, bufferVariance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } @@ -60,8 +58,7 @@ __kernel void kernel_ocl_filter_get_feature(int sample, ccl_global float *variance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - char use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + get_global_id(0); int y = prefilter_rect.y + get_global_id(1); @@ -73,8 +70,7 @@ __kernel void kernel_ocl_filter_get_feature(int sample, mean, variance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl index b7108f3d0f8..66b6e19de84 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel.cl @@ -50,7 +50,6 @@ __kernel void kernel_ocl_path_trace( ccl_constant KernelData *data, ccl_global float *buffer, - ccl_global uint *rng_state, KERNEL_BUFFER_PARAMS, @@ -68,16 +67,15 @@ __kernel void kernel_ocl_path_trace( int y = sy + ccl_global_id(1); if(x < sx + sw && y < sy + sh) - kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride); + kernel_path_trace(kg, buffer, sample, x, y, offset, stride); } #else /* __COMPILE_ONLY_MEGAKERNEL__ */ -__kernel void kernel_ocl_shader( +__kernel void kernel_ocl_displace( ccl_constant KernelData *data, ccl_global uint4 *input, ccl_global float4 *output, - ccl_global float *output_luma, KERNEL_BUFFER_PARAMS, @@ -93,13 +91,29 @@ __kernel void kernel_ocl_shader( int x = sx + ccl_global_id(0); if(x < sx + sw) { - kernel_shader_evaluate(kg, - input, - output, - output_luma, - (ShaderEvalType)type, - x, - sample); + kernel_displace_evaluate(kg, input, output, x); + } +} +__kernel void kernel_ocl_background( + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, + + KERNEL_BUFFER_PARAMS, + + int type, int sx, int sw, int offset, int sample) +{ + KernelGlobals kglobals, *kg = &kglobals; + + kg->data = data; + + kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); + kernel_set_buffer_info(kg); + + int x = sx + ccl_global_id(0); + + if(x < sx + sw) { + kernel_background_evaluate(kg, input, output, x); } } diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl index 95b35e40a45..7125348a49f 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl @@ -24,7 +24,6 @@ __kernel void kernel_ocl_path_trace_data_init( ccl_global void *split_data_buffer, int num_elements, ccl_global char *ray_state, - ccl_global uint *rng_state, KERNEL_BUFFER_PARAMS, int start_sample, int end_sample, @@ -41,7 +40,6 @@ __kernel void kernel_ocl_path_trace_data_init( split_data_buffer, num_elements, ray_state, - rng_state, KERNEL_BUFFER_ARGS, start_sample, end_sample, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h index 591c3846ef2..6aa7681cbed 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h +++ b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h @@ -23,7 +23,6 @@ __kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)( ccl_global void *split_data_buffer, ccl_global char *ray_state, - ccl_global uint *rng_state, KERNEL_BUFFER_PARAMS, @@ -42,11 +41,10 @@ __kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)( 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; + kernel_split_params.tile.buffer = buffer; split_data_init(kg, &kernel_split_state, ccl_global_size(0)*ccl_global_size(1), split_data_buffer, ray_state); diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp index 9a37e0987aa..6b3a996ca12 100644 --- a/intern/cycles/kernel/osl/osl_shader.cpp +++ b/intern/cycles/kernel/osl/osl_shader.cpp @@ -348,14 +348,12 @@ void OSLShader::eval_volume(KernelGlobals *kg, ShaderData *sd, PathState *state, /* Displacement */ -void OSLShader::eval_displacement(KernelGlobals *kg, ShaderData *sd) +void OSLShader::eval_displacement(KernelGlobals *kg, ShaderData *sd, PathState *state) { /* setup shader globals from shader data */ OSLThreadData *tdata = kg->osl_tdata; - PathState state = {0}; - - shaderdata_to_shaderglobals(kg, sd, &state, 0, tdata); + shaderdata_to_shaderglobals(kg, sd, state, 0, tdata); /* execute shader */ OSL::ShadingSystem *ss = (OSL::ShadingSystem*)kg->osl_ss; diff --git a/intern/cycles/kernel/osl/osl_shader.h b/intern/cycles/kernel/osl/osl_shader.h index f7020d1223d..6b392b25cf7 100644 --- a/intern/cycles/kernel/osl/osl_shader.h +++ b/intern/cycles/kernel/osl/osl_shader.h @@ -56,7 +56,7 @@ public: static void eval_surface(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag); static void eval_background(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag); static void eval_volume(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag); - static void eval_displacement(KernelGlobals *kg, ShaderData *sd); + static void eval_displacement(KernelGlobals *kg, ShaderData *sd, PathState *state); /* attributes */ static int find_attribute(KernelGlobals *kg, const ShaderData *sd, uint id, AttributeDescriptor *desc); diff --git a/intern/cycles/kernel/split/kernel_buffer_update.h b/intern/cycles/kernel/split/kernel_buffer_update.h index c9e7deddafa..511334e0550 100644 --- a/intern/cycles/kernel/split/kernel_buffer_update.h +++ b/intern/cycles/kernel/split/kernel_buffer_update.h @@ -75,8 +75,6 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, if(ray_index != QUEUE_EMPTY_SLOT) { #endif - int stride = kernel_split_params.stride; - ccl_global char *ray_state = kernel_split_state.ray_state; ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; @@ -86,7 +84,7 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, if(IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) { uint sample = state->sample; uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.buffer + buffer_offset; + ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; /* accumulate result in output buffer */ kernel_write_result(kg, buffer, sample, L); @@ -96,27 +94,27 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, if(IS_STATE(ray_state, ray_index, RAY_TO_REGENERATE)) { /* We have completed current work; So get next work */ + ccl_global uint *work_pools = kernel_split_params.work_pools; + uint total_work_size = kernel_split_params.total_work_size; uint work_index; - if(!get_next_work(kg, ray_index, &work_index)) { + + if(!get_next_work(kg, work_pools, total_work_size, ray_index, &work_index)) { /* If work is invalid, this means no more work is available and the thread may exit */ ASSIGN_RAY_STATE(ray_state, ray_index, RAY_INACTIVE); } if(IS_STATE(ray_state, ray_index, RAY_TO_REGENERATE)) { + ccl_global WorkTile *tile = &kernel_split_params.tile; uint x, y, sample; - get_work_pixel(kg, work_index, &x, &y, &sample); - - /* Remap rng_state to current pixel. */ - ccl_global uint *rng_state = kernel_split_params.rng_state; - rng_state += kernel_split_params.offset + x + y*stride; + get_work_pixel(tile, work_index, &x, &y, &sample); /* Store buffer offset for writing to passes. */ - uint buffer_offset = (kernel_split_params.offset + x + y*stride) * kernel_data.film.pass_stride; + uint buffer_offset = (tile->offset + x + y*tile->stride) * kernel_data.film.pass_stride; kernel_split_state.buffer_offset[ray_index] = buffer_offset; /* Initialize random numbers and ray. */ uint rng_hash; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, ray); if(ray->t != 0.0f) { /* Initialize throughput, path radiance, Ray, PathState; @@ -132,12 +130,6 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, enqueue_flag = 1; } else { - /* These rays do not participate in path-iteration. */ - float4 L_rad = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - /* Accumulate result in output buffer. */ - ccl_global float *buffer = kernel_split_params.buffer + buffer_offset; - kernel_write_pass_float4(buffer, sample, L_rad); - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_TO_REGENERATE); } } diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h index 2c042dfde6f..77fb61b80a8 100644 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ b/intern/cycles/kernel/split/kernel_data_init.h @@ -23,22 +23,6 @@ CCL_NAMESPACE_BEGIN * The number of elements in the queues is initialized to 0; */ -/* Distributes an amount of work across all threads - * note: work done inside the loop may not show up to all threads till after - * the current kernel has completed - */ -#define parallel_for(kg, iter_name, work_size) \ - for(size_t _size = (work_size), \ - _global_size = ccl_global_size(0) * ccl_global_size(1), \ - _n = _size / _global_size, \ - _thread = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0), \ - iter_name = (_n > 0) ? (_thread * _n) : (_thread) \ - ; \ - (iter_name < (_thread+1) * _n) || (iter_name == _n * _global_size + _thread && _thread < _size % _global_size) \ - ; \ - iter_name = (iter_name != (_thread+1) * _n - 1) ? (iter_name + 1) : (_n * _global_size + _thread) \ - ) - #ifndef __KERNEL_CPU__ ccl_device void kernel_data_init( #else @@ -49,7 +33,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( ccl_global void *split_data_buffer, int num_elements, ccl_global char *ray_state, - ccl_global uint *rng_state, #ifdef __KERNEL_OPENCL__ KERNEL_BUFFER_PARAMS, @@ -73,28 +56,27 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( kg->data = data; #endif - kernel_split_params.x = sx; - kernel_split_params.y = sy; - kernel_split_params.w = sw; - kernel_split_params.h = sh; + kernel_split_params.tile.x = sx; + kernel_split_params.tile.y = sy; + kernel_split_params.tile.w = sw; + kernel_split_params.tile.h = sh; + + kernel_split_params.tile.start_sample = start_sample; + kernel_split_params.tile.num_samples = num_samples; - kernel_split_params.offset = offset; - kernel_split_params.stride = stride; + kernel_split_params.tile.offset = offset; + kernel_split_params.tile.stride = stride; - kernel_split_params.rng_state = rng_state; + kernel_split_params.tile.buffer = buffer; - kernel_split_params.start_sample = start_sample; - kernel_split_params.end_sample = end_sample; + kernel_split_params.total_work_size = sw * sh * num_samples; kernel_split_params.work_pools = work_pools; - kernel_split_params.num_samples = num_samples; kernel_split_params.queue_index = Queue_index; kernel_split_params.queue_size = queuesize; kernel_split_params.use_queues_flag = use_queues_flag; - kernel_split_params.buffer = buffer; - split_data_init(kg, &kernel_split_state, num_elements, split_data_buffer, ray_state); #ifdef __KERNEL_OPENCL__ @@ -121,42 +103,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( */ *use_queues_flag = 0; } - - /* zero the tiles pixels and initialize rng_state if this is the first sample */ - if(start_sample == 0) { - int pass_stride = kernel_data.film.pass_stride; - -#ifdef __KERNEL_CPU__ - for(int y = sy; y < sy + sh; y++) { - int index = offset + y * stride; - memset(buffer + (sx + index) * pass_stride, 0, sizeof(float) * pass_stride * sw); - for(int x = sx; x < sx + sw; x++) { - rng_state[index + x] = hash_int_2d(x, y); - } - } -#else - parallel_for(kg, i, sw * sh * pass_stride) { - int pixel = i / pass_stride; - int pass = i % pass_stride; - - int x = sx + pixel % sw; - int y = sy + pixel / sw; - - int index = (offset + x + y*stride) * pass_stride + pass; - - *(buffer + index) = 0.0f; - } - - parallel_for(kg, i, sw * sh) { - int x = sx + i % sw; - int y = sy + i / sw; - - int index = (offset + x + y*stride); - *(rng_state + index) = hash_int_2d(x, y); - } -#endif - } - #endif /* KERENL_STUB */ } diff --git a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h index dffd291012d..906bad8ceb6 100644 --- a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h +++ b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h @@ -98,7 +98,7 @@ ccl_device void kernel_holdout_emission_blurring_pathtermination_ao( if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.buffer + buffer_offset; + ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; ShaderData *emission_sd = &kernel_split_state.sd_DL_shadow[ray_index]; diff --git a/intern/cycles/kernel/split/kernel_path_init.h b/intern/cycles/kernel/split/kernel_path_init.h index 0ab2289348b..5ad62b585fe 100644 --- a/intern/cycles/kernel/split/kernel_path_init.h +++ b/intern/cycles/kernel/split/kernel_path_init.h @@ -30,29 +30,28 @@ ccl_device void kernel_path_init(KernelGlobals *kg) { kernel_split_state.ray_state[ray_index] = RAY_ACTIVE; /* Get work. */ + ccl_global uint *work_pools = kernel_split_params.work_pools; + uint total_work_size = kernel_split_params.total_work_size; uint work_index; - if(!get_next_work(kg, ray_index, &work_index)) { + + if(!get_next_work(kg, work_pools, total_work_size, ray_index, &work_index)) { /* No more work, mark ray as inactive */ kernel_split_state.ray_state[ray_index] = RAY_INACTIVE; return; } + ccl_global WorkTile *tile = &kernel_split_params.tile; uint x, y, sample; - get_work_pixel(kg, work_index, &x, &y, &sample); - - /* Remap rng_state and buffer to current pixel. */ - ccl_global uint *rng_state = kernel_split_params.rng_state; - rng_state += kernel_split_params.offset + x + y*kernel_split_params.stride; + get_work_pixel(tile, work_index, &x, &y, &sample); /* Store buffer offset for writing to passes. */ - uint buffer_offset = (kernel_split_params.offset + x + y*kernel_split_params.stride) * kernel_data.film.pass_stride; + uint buffer_offset = (tile->offset + x + y*tile->stride) * kernel_data.film.pass_stride; kernel_split_state.buffer_offset[ray_index] = buffer_offset; /* Initialize random numbers and ray. */ uint rng_hash; kernel_path_trace_setup(kg, - rng_state, sample, x, y, &rng_hash, @@ -75,11 +74,6 @@ ccl_device void kernel_path_init(KernelGlobals *kg) { #endif } else { - /* These rays do not participate in path-iteration. */ - float4 L_rad = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - /* Accumulate result in output buffer. */ - ccl_global float *buffer = kernel_split_params.buffer + buffer_offset; - kernel_write_pass_float4(buffer, sample, L_rad); ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_TO_REGENERATE); } } diff --git a/intern/cycles/kernel/split/kernel_split_data_types.h b/intern/cycles/kernel/split/kernel_split_data_types.h index c58c8463f5c..b0e6e5f5250 100644 --- a/intern/cycles/kernel/split/kernel_split_data_types.h +++ b/intern/cycles/kernel/split/kernel_split_data_types.h @@ -22,28 +22,15 @@ CCL_NAMESPACE_BEGIN /* parameters used by the split kernels, we use a single struct to avoid passing these to each kernel */ typedef struct SplitParams { - int x; - int y; - int w; - int h; - - int offset; - int stride; - - ccl_global uint *rng_state; - - int start_sample; - int end_sample; + WorkTile tile; + uint total_work_size; ccl_global unsigned int *work_pools; - unsigned int num_samples; ccl_global int *queue_index; int queue_size; ccl_global char *use_queues_flag; - ccl_global float *buffer; - /* Place for storing sd->flag. AMD GPU OpenCL compiler workaround */ int dummy_sd_flag; } SplitParams; diff --git a/intern/cycles/kernel/svm/svm_blackbody.h b/intern/cycles/kernel/svm/svm_blackbody.h index b750ad87b7f..51590b18505 100644 --- a/intern/cycles/kernel/svm/svm_blackbody.h +++ b/intern/cycles/kernel/svm/svm_blackbody.h @@ -41,8 +41,7 @@ ccl_device void svm_node_blackbody(KernelGlobals *kg, ShaderData *sd, float *sta float3 color_rgb = svm_math_blackbody_color(temperature); - if(stack_valid(col_offset)) - stack_store_float3(stack, col_offset, color_rgb); + stack_store_float3(stack, col_offset, color_rgb); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/svm/svm_math_util.h b/intern/cycles/kernel/svm/svm_math_util.h index a7f15de7325..1ce7777aac3 100644 --- a/intern/cycles/kernel/svm/svm_math_util.h +++ b/intern/cycles/kernel/svm/svm_math_util.h @@ -100,71 +100,64 @@ ccl_device float svm_math(NodeMath type, float Fac1, float Fac2) return Fac; } -ccl_device float3 svm_math_blackbody_color(float t) { - /* Calculate color in range 800..12000 using an approximation - * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B - * Max absolute error for RGB is (0.00095, 0.00077, 0.00057), - * which is enough to get the same 8 bit/channel color. - */ - - const float rc[6][3] = { - { 2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f }, - { 3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f }, - { 4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f }, - { 4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f }, - { 4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f }, - { 3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f }, - }; - - const float gc[6][3] = { - { -7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f }, - { -1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f }, - { -1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f }, - { -1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f }, - { -1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f }, - { -5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f }, - }; - - const float bc[6][4] = { - { 0.0f, 0.0f, 0.0f, 0.0f }, /* zeros should be optimized by compiler */ - { 0.0f, 0.0f, 0.0f, 0.0f }, - { 0.0f, 0.0f, 0.0f, 0.0f }, - { -2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f }, - { -2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f }, - { 6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f }, - }; - - int i; +/* Calculate color in range 800..12000 using an approximation + * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B + * Max absolute error for RGB is (0.00095, 0.00077, 0.00057), + * which is enough to get the same 8 bit/channel color. + */ + +ccl_static_constant float blackbody_table_r[6][3] = { + { 2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f }, + { 3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f }, + { 4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f }, + { 4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f }, + { 4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f }, + { 3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f }, +}; + +ccl_static_constant float blackbody_table_g[6][3] = { + { -7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f }, + { -1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f }, + { -1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f }, + { -1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f }, + { -1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f }, + { -5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f }, +}; + +ccl_static_constant float blackbody_table_b[6][4] = { + { 0.0f, 0.0f, 0.0f, 0.0f }, /* zeros should be optimized by compiler */ + { 0.0f, 0.0f, 0.0f, 0.0f }, + { 0.0f, 0.0f, 0.0f, 0.0f }, + { -2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f }, + { -2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f }, + { 6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f }, +}; + + +ccl_device float3 svm_math_blackbody_color(float t) +{ if(t >= 12000.0f) { return make_float3(0.826270103f, 0.994478524f, 1.56626022f); } - else if(t >= 6365.0f) { - i = 5; - } - else if(t >= 3315.0f) { - i = 4; - } - else if(t >= 1902.0f) { - i = 3; - } - else if(t >= 1449.0f) { - i = 2; - } - else if(t >= 1167.0f) { - i = 1; - } - else if(t >= 965.0f) { - i = 0; - } - else { + else if(t < 965.0f) { /* For 800 <= t < 965 color does not change in OSL implementation, so keep color the same */ return make_float3(4.70366907f, 0.0f, 0.0f); } + int i = (t >= 6365.0f)? 5: + (t >= 3315.0f)? 4: + (t >= 1902.0f)? 3: + (t >= 1449.0f)? 2: + (t >= 1167.0f)? 1: 0; + + ccl_constant float *r = blackbody_table_r[i]; + ccl_constant float *g = blackbody_table_g[i]; + ccl_constant float *b = blackbody_table_b[i]; + const float t_inv = 1.0f / t; - return make_float3(rc[i][0] * t_inv + rc[i][1] * t + rc[i][2], - gc[i][0] * t_inv + gc[i][1] * t + gc[i][2], - ((bc[i][0] * t + bc[i][1]) * t + bc[i][2]) * t + bc[i][3]); + return make_float3(r[0] * t_inv + r[1] * t + r[2], + g[0] * t_inv + g[1] * t + g[2], + ((b[0] * t + b[1]) * t + b[2]) * t + b[3]); } ccl_device_inline float3 svm_math_gamma_color(float3 color, float gamma) diff --git a/intern/cycles/kernel/svm/svm_wavelength.h b/intern/cycles/kernel/svm/svm_wavelength.h index 57030f3979d..855b356b397 100644 --- a/intern/cycles/kernel/svm/svm_wavelength.h +++ b/intern/cycles/kernel/svm/svm_wavelength.h @@ -34,44 +34,44 @@ CCL_NAMESPACE_BEGIN /* Wavelength to RGB */ +// CIE colour matching functions xBar, yBar, and zBar for +// wavelengths from 380 through 780 nanometers, every 5 +// nanometers. For a wavelength lambda in this range: +// cie_colour_match[(lambda - 380) / 5][0] = xBar +// cie_colour_match[(lambda - 380) / 5][1] = yBar +// cie_colour_match[(lambda - 380) / 5][2] = zBar +ccl_static_constant float cie_colour_match[81][3] = { + {0.0014f,0.0000f,0.0065f}, {0.0022f,0.0001f,0.0105f}, {0.0042f,0.0001f,0.0201f}, + {0.0076f,0.0002f,0.0362f}, {0.0143f,0.0004f,0.0679f}, {0.0232f,0.0006f,0.1102f}, + {0.0435f,0.0012f,0.2074f}, {0.0776f,0.0022f,0.3713f}, {0.1344f,0.0040f,0.6456f}, + {0.2148f,0.0073f,1.0391f}, {0.2839f,0.0116f,1.3856f}, {0.3285f,0.0168f,1.6230f}, + {0.3483f,0.0230f,1.7471f}, {0.3481f,0.0298f,1.7826f}, {0.3362f,0.0380f,1.7721f}, + {0.3187f,0.0480f,1.7441f}, {0.2908f,0.0600f,1.6692f}, {0.2511f,0.0739f,1.5281f}, + {0.1954f,0.0910f,1.2876f}, {0.1421f,0.1126f,1.0419f}, {0.0956f,0.1390f,0.8130f}, + {0.0580f,0.1693f,0.6162f}, {0.0320f,0.2080f,0.4652f}, {0.0147f,0.2586f,0.3533f}, + {0.0049f,0.3230f,0.2720f}, {0.0024f,0.4073f,0.2123f}, {0.0093f,0.5030f,0.1582f}, + {0.0291f,0.6082f,0.1117f}, {0.0633f,0.7100f,0.0782f}, {0.1096f,0.7932f,0.0573f}, + {0.1655f,0.8620f,0.0422f}, {0.2257f,0.9149f,0.0298f}, {0.2904f,0.9540f,0.0203f}, + {0.3597f,0.9803f,0.0134f}, {0.4334f,0.9950f,0.0087f}, {0.5121f,1.0000f,0.0057f}, + {0.5945f,0.9950f,0.0039f}, {0.6784f,0.9786f,0.0027f}, {0.7621f,0.9520f,0.0021f}, + {0.8425f,0.9154f,0.0018f}, {0.9163f,0.8700f,0.0017f}, {0.9786f,0.8163f,0.0014f}, + {1.0263f,0.7570f,0.0011f}, {1.0567f,0.6949f,0.0010f}, {1.0622f,0.6310f,0.0008f}, + {1.0456f,0.5668f,0.0006f}, {1.0026f,0.5030f,0.0003f}, {0.9384f,0.4412f,0.0002f}, + {0.8544f,0.3810f,0.0002f}, {0.7514f,0.3210f,0.0001f}, {0.6424f,0.2650f,0.0000f}, + {0.5419f,0.2170f,0.0000f}, {0.4479f,0.1750f,0.0000f}, {0.3608f,0.1382f,0.0000f}, + {0.2835f,0.1070f,0.0000f}, {0.2187f,0.0816f,0.0000f}, {0.1649f,0.0610f,0.0000f}, + {0.1212f,0.0446f,0.0000f}, {0.0874f,0.0320f,0.0000f}, {0.0636f,0.0232f,0.0000f}, + {0.0468f,0.0170f,0.0000f}, {0.0329f,0.0119f,0.0000f}, {0.0227f,0.0082f,0.0000f}, + {0.0158f,0.0057f,0.0000f}, {0.0114f,0.0041f,0.0000f}, {0.0081f,0.0029f,0.0000f}, + {0.0058f,0.0021f,0.0000f}, {0.0041f,0.0015f,0.0000f}, {0.0029f,0.0010f,0.0000f}, + {0.0020f,0.0007f,0.0000f}, {0.0014f,0.0005f,0.0000f}, {0.0010f,0.0004f,0.0000f}, + {0.0007f,0.0002f,0.0000f}, {0.0005f,0.0002f,0.0000f}, {0.0003f,0.0001f,0.0000f}, + {0.0002f,0.0001f,0.0000f}, {0.0002f,0.0001f,0.0000f}, {0.0001f,0.0000f,0.0000f}, + {0.0001f,0.0000f,0.0000f}, {0.0001f,0.0000f,0.0000f}, {0.0000f,0.0000f,0.0000f} +}; + ccl_device void svm_node_wavelength(ShaderData *sd, float *stack, uint wavelength, uint color_out) { - // CIE colour matching functions xBar, yBar, and zBar for - // wavelengths from 380 through 780 nanometers, every 5 - // nanometers. For a wavelength lambda in this range: - // cie_colour_match[(lambda - 380) / 5][0] = xBar - // cie_colour_match[(lambda - 380) / 5][1] = yBar - // cie_colour_match[(lambda - 380) / 5][2] = zBar - const float cie_colour_match[81][3] = { - {0.0014f,0.0000f,0.0065f}, {0.0022f,0.0001f,0.0105f}, {0.0042f,0.0001f,0.0201f}, - {0.0076f,0.0002f,0.0362f}, {0.0143f,0.0004f,0.0679f}, {0.0232f,0.0006f,0.1102f}, - {0.0435f,0.0012f,0.2074f}, {0.0776f,0.0022f,0.3713f}, {0.1344f,0.0040f,0.6456f}, - {0.2148f,0.0073f,1.0391f}, {0.2839f,0.0116f,1.3856f}, {0.3285f,0.0168f,1.6230f}, - {0.3483f,0.0230f,1.7471f}, {0.3481f,0.0298f,1.7826f}, {0.3362f,0.0380f,1.7721f}, - {0.3187f,0.0480f,1.7441f}, {0.2908f,0.0600f,1.6692f}, {0.2511f,0.0739f,1.5281f}, - {0.1954f,0.0910f,1.2876f}, {0.1421f,0.1126f,1.0419f}, {0.0956f,0.1390f,0.8130f}, - {0.0580f,0.1693f,0.6162f}, {0.0320f,0.2080f,0.4652f}, {0.0147f,0.2586f,0.3533f}, - {0.0049f,0.3230f,0.2720f}, {0.0024f,0.4073f,0.2123f}, {0.0093f,0.5030f,0.1582f}, - {0.0291f,0.6082f,0.1117f}, {0.0633f,0.7100f,0.0782f}, {0.1096f,0.7932f,0.0573f}, - {0.1655f,0.8620f,0.0422f}, {0.2257f,0.9149f,0.0298f}, {0.2904f,0.9540f,0.0203f}, - {0.3597f,0.9803f,0.0134f}, {0.4334f,0.9950f,0.0087f}, {0.5121f,1.0000f,0.0057f}, - {0.5945f,0.9950f,0.0039f}, {0.6784f,0.9786f,0.0027f}, {0.7621f,0.9520f,0.0021f}, - {0.8425f,0.9154f,0.0018f}, {0.9163f,0.8700f,0.0017f}, {0.9786f,0.8163f,0.0014f}, - {1.0263f,0.7570f,0.0011f}, {1.0567f,0.6949f,0.0010f}, {1.0622f,0.6310f,0.0008f}, - {1.0456f,0.5668f,0.0006f}, {1.0026f,0.5030f,0.0003f}, {0.9384f,0.4412f,0.0002f}, - {0.8544f,0.3810f,0.0002f}, {0.7514f,0.3210f,0.0001f}, {0.6424f,0.2650f,0.0000f}, - {0.5419f,0.2170f,0.0000f}, {0.4479f,0.1750f,0.0000f}, {0.3608f,0.1382f,0.0000f}, - {0.2835f,0.1070f,0.0000f}, {0.2187f,0.0816f,0.0000f}, {0.1649f,0.0610f,0.0000f}, - {0.1212f,0.0446f,0.0000f}, {0.0874f,0.0320f,0.0000f}, {0.0636f,0.0232f,0.0000f}, - {0.0468f,0.0170f,0.0000f}, {0.0329f,0.0119f,0.0000f}, {0.0227f,0.0082f,0.0000f}, - {0.0158f,0.0057f,0.0000f}, {0.0114f,0.0041f,0.0000f}, {0.0081f,0.0029f,0.0000f}, - {0.0058f,0.0021f,0.0000f}, {0.0041f,0.0015f,0.0000f}, {0.0029f,0.0010f,0.0000f}, - {0.0020f,0.0007f,0.0000f}, {0.0014f,0.0005f,0.0000f}, {0.0010f,0.0004f,0.0000f}, - {0.0007f,0.0002f,0.0000f}, {0.0005f,0.0002f,0.0000f}, {0.0003f,0.0001f,0.0000f}, - {0.0002f,0.0001f,0.0000f}, {0.0002f,0.0001f,0.0000f}, {0.0001f,0.0000f,0.0000f}, - {0.0001f,0.0000f,0.0000f}, {0.0001f,0.0000f,0.0000f}, {0.0000f,0.0000f,0.0000f} - }; - float lambda_nm = stack_load_float(stack, wavelength); float ii = (lambda_nm-380.0f) * (1.0f/5.0f); // scaled 0..80 int i = float_to_int(ii); @@ -82,7 +82,7 @@ ccl_device void svm_node_wavelength(ShaderData *sd, float *stack, uint wavelengt } else { ii -= i; - const float *c = cie_colour_match[i]; + ccl_constant float *c = cie_colour_match[i]; color = interp(make_float3(c[0], c[1], c[2]), make_float3(c[3], c[4], c[5]), ii); } @@ -92,8 +92,7 @@ ccl_device void svm_node_wavelength(ShaderData *sd, float *stack, uint wavelengt /* Clamp to zero if values are smaller */ color = max(color, make_float3(0.0f, 0.0f, 0.0f)); - if(stack_valid(color_out)) - stack_store_float3(stack, color_out, color); + stack_store_float3(stack, color_out, color); } CCL_NAMESPACE_END diff --git a/intern/cycles/render/bake.cpp b/intern/cycles/render/bake.cpp index c0fcd517390..2bedf3668f7 100644 --- a/intern/cycles/render/bake.cpp +++ b/intern/cycles/render/bake.cpp @@ -174,6 +174,7 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre device->mem_alloc("bake_input", d_input, MEM_READ_ONLY); device->mem_copy_to(d_input); device->mem_alloc("bake_output", d_output, MEM_READ_WRITE); + device->mem_zero(d_output); DeviceTask task(DeviceTask::SHADER); task.shader_input = d_input.device_pointer; diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp index ff6f40c010e..a7a1260e425 100644 --- a/intern/cycles/render/buffers.cpp +++ b/intern/cycles/render/buffers.cpp @@ -108,7 +108,6 @@ RenderTile::RenderTile() stride = 0; buffer = 0; - rng_state = 0; buffers = NULL; } @@ -131,11 +130,6 @@ void RenderBuffers::device_free() device->mem_free(buffer); buffer.clear(); } - - if(rng_state.device_pointer) { - device->mem_free(rng_state); - rng_state.clear(); - } } void RenderBuffers::reset(Device *device, BufferParams& params_) @@ -149,11 +143,13 @@ void RenderBuffers::reset(Device *device, BufferParams& params_) buffer.resize(params.width*params.height*params.get_passes_size()); device->mem_alloc("render_buffer", buffer, MEM_READ_WRITE); device->mem_zero(buffer); +} - /* allocate rng state */ - rng_state.resize(params.width, params.height); - - device->mem_alloc("rng_state", rng_state, MEM_READ_WRITE); +void RenderBuffers::zero(Device *device) +{ + if(buffer.device_pointer) { + device->mem_zero(buffer); + } } bool RenderBuffers::copy_from_device(Device *from_device) diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h index e56556c8abe..2780fc8a68d 100644 --- a/intern/cycles/render/buffers.h +++ b/intern/cycles/render/buffers.h @@ -74,8 +74,6 @@ public: /* float buffer */ device_vector<float> buffer; - /* random number generator state */ - device_vector<uint> rng_state; Device *device; @@ -83,6 +81,7 @@ public: ~RenderBuffers(); void reset(Device *device, BufferParams& params); + void zero(Device *device); bool copy_from_device(Device *from_device = NULL); bool get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels); @@ -149,7 +148,6 @@ public: int tile_index; device_ptr buffer; - device_ptr rng_state; RenderBuffers *buffers; diff --git a/intern/cycles/render/light.cpp b/intern/cycles/render/light.cpp index 4adc00bc839..6a7f985b756 100644 --- a/intern/cycles/render/light.cpp +++ b/intern/cycles/render/light.cpp @@ -60,6 +60,7 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res device->mem_alloc("shade_background_pixels_input", d_input, MEM_READ_ONLY); device->mem_copy_to(d_input); device->mem_alloc("shade_background_pixels_output", d_output, MEM_WRITE_ONLY); + device->mem_zero(d_output); DeviceTask main_task(DeviceTask::SHADER); main_task.shader_input = d_input.device_pointer; diff --git a/intern/cycles/render/mesh_displace.cpp b/intern/cycles/render/mesh_displace.cpp index 4ca20cf7ef3..350a56bf185 100644 --- a/intern/cycles/render/mesh_displace.cpp +++ b/intern/cycles/render/mesh_displace.cpp @@ -124,6 +124,7 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me device->mem_alloc("displace_input", d_input, MEM_READ_ONLY); device->mem_copy_to(d_input); device->mem_alloc("displace_output", d_output, MEM_WRITE_ONLY); + device->mem_zero(d_output); DeviceTask task(DeviceTask::SHADER); task.shader_input = d_input.device_pointer; diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp index f68efe38add..f1ff6b49b71 100644 --- a/intern/cycles/render/session.cpp +++ b/intern/cycles/render/session.cpp @@ -384,7 +384,6 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile) tile_manager.state.buffer.get_offset_stride(rtile.offset, rtile.stride); rtile.buffer = buffers->buffer.device_pointer; - rtile.rng_state = buffers->rng_state.device_pointer; rtile.buffers = buffers; tile->buffers = buffers; @@ -442,7 +441,6 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile) tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride); rtile.buffer = tile->buffers->buffer.device_pointer; - rtile.rng_state = tile->buffers->rng_state.device_pointer; rtile.buffers = tile->buffers; rtile.sample = 0; @@ -972,7 +970,12 @@ void Session::update_status_time(bool show_pause, bool show_done) void Session::render() { - /* add path trace task */ + /* Clear buffers. */ + if(buffers && tile_manager.state.sample == 0) { + buffers->zero(device); + } + + /* Add path trace task. */ DeviceTask task(DeviceTask::RENDER); task.acquire_tile = function_bind(&Session::acquire_tile, this, _1, _2); diff --git a/intern/cycles/util/util_defines.h b/intern/cycles/util/util_defines.h index ae654092c87..98944a19022 100644 --- a/intern/cycles/util/util_defines.h +++ b/intern/cycles/util/util_defines.h @@ -30,7 +30,8 @@ # define ccl_device static inline # define ccl_device_noinline static # define ccl_global -# define ccl_constant +# define ccl_static_constant static const +# define ccl_constant const # define ccl_local # define ccl_local_param # define ccl_private diff --git a/intern/ghost/intern/GHOST_WindowWin32.cpp b/intern/ghost/intern/GHOST_WindowWin32.cpp index fc4979291ec..ee0d7ef460c 100644 --- a/intern/ghost/intern/GHOST_WindowWin32.cpp +++ b/intern/ghost/intern/GHOST_WindowWin32.cpp @@ -402,21 +402,29 @@ void GHOST_WindowWin32::getClientBounds(GHOST_Rect &bounds) const { RECT rect; POINT coord; - ::GetClientRect(m_hWnd, &rect); + if (!IsIconic(m_hWnd)) { + ::GetClientRect(m_hWnd, &rect); - coord.x = rect.left; - coord.y = rect.top; - ::ClientToScreen(m_hWnd, &coord); + coord.x = rect.left; + coord.y = rect.top; + ::ClientToScreen(m_hWnd, &coord); - bounds.m_l = coord.x; - bounds.m_t = coord.y; + bounds.m_l = coord.x; + bounds.m_t = coord.y; - coord.x = rect.right; - coord.y = rect.bottom; - ::ClientToScreen(m_hWnd, &coord); + coord.x = rect.right; + coord.y = rect.bottom; + ::ClientToScreen(m_hWnd, &coord); - bounds.m_r = coord.x; - bounds.m_b = coord.y; + bounds.m_r = coord.x; + bounds.m_b = coord.y; + } + else { + bounds.m_b = 0; + bounds.m_l = 0; + bounds.m_r = 0; + bounds.m_t = 0; + } } diff --git a/release/scripts/modules/bl_previews_utils/bl_previews_render.py b/release/scripts/modules/bl_previews_utils/bl_previews_render.py index b4f0b8dd3b8..71208ef3485 100644 --- a/release/scripts/modules/bl_previews_utils/bl_previews_render.py +++ b/release/scripts/modules/bl_previews_utils/bl_previews_render.py @@ -278,7 +278,7 @@ def do_previews(do_objects, do_groups, do_scenes, do_data_intern): bbox[1].z = v.z def objects_bbox_calc(camera, objects, offset_matrix): - bbox = (Vector((1e9, 1e9, 1e9)), Vector((-1e9, -1e9, -1e9))) + bbox = (Vector((1e24, 1e24, 1e24)), Vector((-1e24, -1e24, -1e24))) for obname, libpath in objects: ob = bpy.data.objects[obname, libpath] object_bbox_merge(bbox, ob, camera, offset_matrix) @@ -305,6 +305,17 @@ def do_previews(do_objects, do_groups, do_scenes, do_data_intern): cos = objects_bbox_calc(camera, objects, offset_matrix) loc, ortho_scale = camera.camera_fit_coords(scene, cos) camera.location = loc + # Set camera clipping accordingly to computed bbox. + min_dist = 1e24 + max_dist = -1e24 + for co in zip(*(iter(cos),) * 3): + dist = (Vector(co) - loc).length + if dist < min_dist: + min_dist = dist + if dist > max_dist: + max_dist = dist + camera.data.clip_start = min_dist / 2 + camera.data.clip_end = max_dist * 2 if lamp: loc, ortho_scale = lamp.camera_fit_coords(scene, cos) lamp.location = loc diff --git a/release/scripts/startup/bl_ui/space_view3d_toolbar.py b/release/scripts/startup/bl_ui/space_view3d_toolbar.py index 0f9dd117c14..5430b108642 100644 --- a/release/scripts/startup/bl_ui/space_view3d_toolbar.py +++ b/release/scripts/startup/bl_ui/space_view3d_toolbar.py @@ -1080,8 +1080,8 @@ class VIEW3D_PT_tools_brush(Panel, View3DPaintPanel): # use_frontface col.separator() - row = col.row() - row.prop(brush, "use_frontface", text="Front Faces Only") + col.prop(brush, "use_frontface", text="Front Faces Only") + col.prop(brush, "use_projected") # direction col.separator() @@ -1131,12 +1131,22 @@ class VIEW3D_PT_tools_brush(Panel, View3DPaintPanel): self.prop_unified_strength(row, context, brush, "strength", text="Strength") self.prop_unified_strength(row, context, brush, "use_pressure_strength") + col.separator() col.prop(brush, "vertex_tool", text="Blend") if brush.vertex_tool != 'SMEAR': col.prop(brush, "use_accumulate") col.separator() + col.prop(brush, "use_frontface", text="Front Faces Only") + row = col.row() + row.prop(brush, "use_frontface_falloff", text="Falloff Angle") + sub = row.row() + sub.active = brush.use_frontface_falloff + sub.prop(brush, "falloff_angle", text="") + + col.prop(brush, "use_projected") + col = layout.column() col.prop(toolsettings, "use_auto_normalize", text="Auto Normalize") col.prop(toolsettings, "use_multipaint", text="Multi-Paint") @@ -1162,15 +1172,23 @@ class VIEW3D_PT_tools_brush(Panel, View3DPaintPanel): self.prop_unified_strength(row, context, brush, "strength", text="Strength") self.prop_unified_strength(row, context, brush, "use_pressure_strength") - # XXX - TODO - # row = col.row(align=True) - # row.prop(brush, "jitter", slider=True) - # row.prop(brush, "use_pressure_jitter", toggle=True, text="") col.separator() col.prop(brush, "vertex_tool", text="Blend") - col.prop(brush, "use_accumulate") col.prop(brush, "use_alpha") + if brush.vertex_tool != 'SMEAR': + col.prop(brush, "use_accumulate") + col.separator() + + col.prop(brush, "use_frontface", text="Front Faces Only") + row = col.row() + row.prop(brush, "use_frontface_falloff", text="Falloff Angle") + sub = row.row() + sub.active = brush.use_frontface_falloff + sub.prop(brush, "falloff_angle", text="") + + col.prop(brush, "use_projected") + col.separator() col.template_ID(settings, "palette", new="palette.new") @@ -1761,19 +1779,7 @@ class VIEW3D_PT_tools_weightpaint_options(Panel, View3DPaintPanel): wpaint = tool_settings.weight_paint col = layout.column() - col.label("Falloff:") - row = col.row() - row.prop(wpaint, "falloff_shape", expand=True) - row = col.row() - row.prop(wpaint, "use_backface_culling") - row = col.row() - row.prop(wpaint, "use_normal_falloff") - sub = row.row() - sub.active = (wpaint.use_normal_falloff) - sub.prop(wpaint, "normal_angle", text="") - col = layout.column() - row = col.row() - row.prop(wpaint, "use_group_restrict") + col.prop(wpaint, "use_group_restrict") obj = context.weight_paint_object if obj.type == 'MESH': @@ -1806,10 +1812,6 @@ class VIEW3D_PT_tools_vertexpaint(Panel, View3DPaintPanel): col = layout.column() col.label("Falloff:") row = col.row() - row.prop(vpaint, "falloff_shape", expand=True) - row = col.row() - row.prop(vpaint, "use_backface_culling") - row = col.row() row.prop(vpaint, "use_normal_falloff") sub = row.row() sub.active = (vpaint.use_normal_falloff) diff --git a/source/blender/blenkernel/BKE_brush.h b/source/blender/blenkernel/BKE_brush.h index c6d4217780e..7c2873046d5 100644 --- a/source/blender/blenkernel/BKE_brush.h +++ b/source/blender/blenkernel/BKE_brush.h @@ -103,9 +103,11 @@ void BKE_brush_alpha_set(struct Scene *scene, struct Brush *brush, float alpha); float BKE_brush_weight_get(const struct Scene *scene, const struct Brush *brush); void BKE_brush_weight_set(const struct Scene *scene, struct Brush *brush, float value); -int BKE_brush_use_locked_size(const struct Scene *scene, const struct Brush *brush); -int BKE_brush_use_alpha_pressure(const struct Scene *scene, const struct Brush *brush); -int BKE_brush_use_size_pressure(const struct Scene *scene, const struct Brush *brush); +bool BKE_brush_use_locked_size(const struct Scene *scene, const struct Brush *brush); +bool BKE_brush_use_alpha_pressure(const struct Scene *scene, const struct Brush *brush); +bool BKE_brush_use_size_pressure(const struct Scene *scene, const struct Brush *brush); + +bool BKE_brush_sculpt_has_secondary_color(const struct Brush *brush); /* scale unprojected radius to reflect a change in the brush's 2D size */ void BKE_brush_scale_unprojected_radius( diff --git a/source/blender/blenkernel/BKE_pbvh.h b/source/blender/blenkernel/BKE_pbvh.h index 8412127c701..9aea5dc95a0 100644 --- a/source/blender/blenkernel/BKE_pbvh.h +++ b/source/blender/blenkernel/BKE_pbvh.h @@ -60,6 +60,8 @@ typedef bool (*BKE_pbvh_SearchCallback)(PBVHNode *node, void *data); typedef void (*BKE_pbvh_HitCallback)(PBVHNode *node, void *data); typedef void (*BKE_pbvh_HitOccludedCallback)(PBVHNode *node, void *data, float *tmin); +typedef void (*BKE_pbvh_SearchNearestCallback)(PBVHNode *node, void *data, float *tmin); + /* Building */ PBVH *BKE_pbvh_new(void); @@ -102,12 +104,12 @@ void BKE_pbvh_raycast( bool BKE_pbvh_node_raycast( PBVH *bvh, PBVHNode *node, float (*origco)[3], bool use_origco, const float ray_start[3], const float ray_normal[3], - float *dist); + float *depth); bool BKE_pbvh_bmesh_node_raycast_detail( PBVHNode *node, const float ray_start[3], const float ray_normal[3], - float *dist, float *r_detail); + float *depth, float *r_detail); /* for orthographic cameras, project the far away ray segment points to the root node so * we can have better precision. */ @@ -115,6 +117,16 @@ void BKE_pbvh_raycast_project_ray_root( PBVH *bvh, bool original, float ray_start[3], float ray_end[3], float ray_normal[3]); +void BKE_pbvh_find_nearest_to_ray( + PBVH *bvh, BKE_pbvh_HitOccludedCallback cb, void *data, + const float ray_start[3], const float ray_normal[3], + bool original); + +bool BKE_pbvh_node_find_nearest_to_ray( + PBVH *bvh, PBVHNode *node, float (*origco)[3], bool use_origco, + const float ray_start[3], const float ray_normal[3], + float *depth, float *dist_sq); + /* Drawing */ void BKE_pbvh_node_draw(PBVHNode *node, void *data); @@ -160,7 +172,7 @@ typedef enum { bool BKE_pbvh_bmesh_update_topology( PBVH *bvh, PBVHTopologyUpdateMode mode, const float center[3], const float view_normal[3], - float radius); + float radius, const bool use_frontface, const bool use_projected); /* Node Access */ diff --git a/source/blender/blenkernel/intern/brush.c b/source/blender/blenkernel/intern/brush.c index 4a54ab26373..aeaead578a1 100644 --- a/source/blender/blenkernel/intern/brush.c +++ b/source/blender/blenkernel/intern/brush.c @@ -821,7 +821,7 @@ int BKE_brush_size_get(const Scene *scene, const Brush *brush) return size; } -int BKE_brush_use_locked_size(const Scene *scene, const Brush *brush) +bool BKE_brush_use_locked_size(const Scene *scene, const Brush *brush) { const short us_flag = scene->toolsettings->unified_paint_settings.flag; @@ -830,7 +830,7 @@ int BKE_brush_use_locked_size(const Scene *scene, const Brush *brush) (brush->flag & BRUSH_LOCK_SIZE); } -int BKE_brush_use_size_pressure(const Scene *scene, const Brush *brush) +bool BKE_brush_use_size_pressure(const Scene *scene, const Brush *brush) { const short us_flag = scene->toolsettings->unified_paint_settings.flag; @@ -839,7 +839,7 @@ int BKE_brush_use_size_pressure(const Scene *scene, const Brush *brush) (brush->flag & BRUSH_SIZE_PRESSURE); } -int BKE_brush_use_alpha_pressure(const Scene *scene, const Brush *brush) +bool BKE_brush_use_alpha_pressure(const Scene *scene, const Brush *brush) { const short us_flag = scene->toolsettings->unified_paint_settings.flag; @@ -848,6 +848,16 @@ int BKE_brush_use_alpha_pressure(const Scene *scene, const Brush *brush) (brush->flag & BRUSH_ALPHA_PRESSURE); } +bool BKE_brush_sculpt_has_secondary_color(const Brush *brush) +{ + return ELEM( + brush->sculpt_tool, SCULPT_TOOL_BLOB, SCULPT_TOOL_DRAW, + SCULPT_TOOL_INFLATE, SCULPT_TOOL_CLAY, SCULPT_TOOL_CLAY_STRIPS, + SCULPT_TOOL_PINCH, SCULPT_TOOL_CREASE, SCULPT_TOOL_LAYER, + SCULPT_TOOL_FLATTEN, SCULPT_TOOL_FILL, SCULPT_TOOL_SCRAPE, + SCULPT_TOOL_MASK); +} + void BKE_brush_unprojected_radius_set(Scene *scene, Brush *brush, float unprojected_radius) { UnifiedPaintSettings *ups = &scene->toolsettings->unified_paint_settings; diff --git a/source/blender/blenkernel/intern/colortools.c b/source/blender/blenkernel/intern/colortools.c index ee0f904c3a6..310255a15c1 100644 --- a/source/blender/blenkernel/intern/colortools.c +++ b/source/blender/blenkernel/intern/colortools.c @@ -1380,7 +1380,7 @@ void scopes_update(Scopes *scopes, ImBuf *ibuf, const ColorManagedViewSettings * /* Keep number of threads in sync with the merge parts below. */ ScopesUpdateData data = { - .scopes = scopes, . ibuf = ibuf, + .scopes = scopes, .ibuf = ibuf, .cm_processor = cm_processor, .display_buffer = display_buffer, .ycc_mode = ycc_mode, .bin_lum = bin_lum, .bin_r = bin_r, .bin_g = bin_g, .bin_b = bin_b, .bin_a = bin_a, }; diff --git a/source/blender/blenkernel/intern/image.c b/source/blender/blenkernel/intern/image.c index f826f655309..33a690b32c6 100644 --- a/source/blender/blenkernel/intern/image.c +++ b/source/blender/blenkernel/intern/image.c @@ -1455,8 +1455,7 @@ static bool do_add_image_extension(char *string, const char imtype, const ImageF if (extension) { /* prefer this in many cases to avoid .png.tga, but in certain cases it breaks */ /* remove any other known image extension */ - if (BLI_testextensie_array(string, imb_ext_image)) - { + if (BLI_testextensie_array(string, imb_ext_image)) { return BLI_replace_extension(string, FILE_MAX, extension); } else { diff --git a/source/blender/blenkernel/intern/mball_tessellate.c b/source/blender/blenkernel/intern/mball_tessellate.c index 1d9c580d45b..d8e3df239df 100644 --- a/source/blender/blenkernel/intern/mball_tessellate.c +++ b/source/blender/blenkernel/intern/mball_tessellate.c @@ -424,13 +424,13 @@ static void make_face(PROCESS *process, int i1, int i2, int i3, int i4) #ifdef USE_ACCUM_NORMAL if (i4 == 0) { normal_tri_v3(n, process->co[i1], process->co[i2], process->co[i3]); - accumulate_vertex_normals( + accumulate_vertex_normals_v3( process->no[i1], process->no[i2], process->no[i3], NULL, n, process->co[i1], process->co[i2], process->co[i3], NULL); } else { normal_quad_v3(n, process->co[i1], process->co[i2], process->co[i3], process->co[i4]); - accumulate_vertex_normals( + accumulate_vertex_normals_v3( process->no[i1], process->no[i2], process->no[i3], process->no[i4], n, process->co[i1], process->co[i2], process->co[i3], process->co[i4]); } diff --git a/source/blender/blenkernel/intern/mesh_evaluate.c b/source/blender/blenkernel/intern/mesh_evaluate.c index 654a25cbb3d..4a1952c798b 100644 --- a/source/blender/blenkernel/intern/mesh_evaluate.c +++ b/source/blender/blenkernel/intern/mesh_evaluate.c @@ -225,7 +225,7 @@ static void mesh_calc_normals_poly_accum_task_cb(void *userdata, const int pidx) } /* accumulate angle weighted face normal */ - /* inline version of #accumulate_vertex_normals_poly */ + /* inline version of #accumulate_vertex_normals_poly_v3 */ { const float *prev_edge = edgevecbuf[nverts - 1]; @@ -334,8 +334,9 @@ void BKE_mesh_calc_normals_tessface( else normal_tri_v3(f_no, mverts[mf->v1].co, mverts[mf->v2].co, mverts[mf->v3].co); - accumulate_vertex_normals(tnorms[mf->v1], tnorms[mf->v2], tnorms[mf->v3], n4, - f_no, mverts[mf->v1].co, mverts[mf->v2].co, mverts[mf->v3].co, c4); + accumulate_vertex_normals_v3( + tnorms[mf->v1], tnorms[mf->v2], tnorms[mf->v3], n4, + f_no, mverts[mf->v1].co, mverts[mf->v2].co, mverts[mf->v3].co, c4); } /* following Mesh convention; we use vertex coordinate itself for normal in this case */ @@ -379,7 +380,7 @@ void BKE_mesh_calc_normals_looptri( f_no, mverts[vtri[0]].co, mverts[vtri[1]].co, mverts[vtri[2]].co); - accumulate_vertex_normals_tri( + accumulate_vertex_normals_tri_v3( tnorms[vtri[0]], tnorms[vtri[1]], tnorms[vtri[2]], f_no, mverts[vtri[0]].co, mverts[vtri[1]].co, mverts[vtri[2]].co); } @@ -845,7 +846,7 @@ static void split_loop_nor_fan_do(LoopSplitTaskDataCommon *common_data, LoopSpli // printf("\thandling edge %d / loop %d\n", mlfan_curr->e, mlfan_curr_index); { - /* Code similar to accumulate_vertex_normals_poly. */ + /* Code similar to accumulate_vertex_normals_poly_v3. */ /* Calculate angle between the two poly edges incident on this vertex. */ const float fac = saacos(dot_v3v3(vec_curr, vec_prev)); /* Accumulate */ diff --git a/source/blender/blenkernel/intern/pbvh.c b/source/blender/blenkernel/intern/pbvh.c index a90a5e2a10b..6fc89eb778a 100644 --- a/source/blender/blenkernel/intern/pbvh.c +++ b/source/blender/blenkernel/intern/pbvh.c @@ -608,7 +608,8 @@ void BKE_pbvh_build_grids(PBVH *bvh, CCGElem **grids, MEM_freeN(prim_bbc); } -void BKE_pbvh_set_ccgdm(PBVH *bvh, CCGDerivedMesh *ccgdm) { +void BKE_pbvh_set_ccgdm(PBVH *bvh, CCGDerivedMesh *ccgdm) +{ bvh->ccgdm = ccgdm; } @@ -1336,7 +1337,8 @@ void BKE_pbvh_get_grid_key(const PBVH *bvh, CCGKey *key) *key = bvh->gridkey; } -CCGDerivedMesh *BKE_pbvh_get_ccgdm(const PBVH *bvh) { +CCGDerivedMesh *BKE_pbvh_get_ccgdm(const PBVH *bvh) +{ return bvh->ccgdm; } @@ -1539,14 +1541,16 @@ void BKE_pbvh_raycast( bool ray_face_intersection_quad( const float ray_start[3], const float ray_normal[3], const float t0[3], const float t1[3], const float t2[3], const float t3[3], - float *dist) + float *depth) { - float dist_test; + float depth_test; - if ((isect_ray_tri_epsilon_v3(ray_start, ray_normal, t0, t1, t2, &dist_test, NULL, 0.1f) && (dist_test < *dist)) || - (isect_ray_tri_epsilon_v3(ray_start, ray_normal, t0, t2, t3, &dist_test, NULL, 0.1f) && (dist_test < *dist))) + if ((isect_ray_tri_epsilon_v3( + ray_start, ray_normal, t0, t1, t2, &depth_test, NULL, 0.1f) && (depth_test < *depth)) || + (isect_ray_tri_epsilon_v3( + ray_start, ray_normal, t0, t2, t3, &depth_test, NULL, 0.1f) && (depth_test < *depth))) { - *dist = dist_test; + *depth = depth_test; return true; } else { @@ -1557,12 +1561,82 @@ bool ray_face_intersection_quad( bool ray_face_intersection_tri( const float ray_start[3], const float ray_normal[3], const float t0[3], const float t1[3], const float t2[3], - float *dist) + float *depth) +{ + float depth_test; + + if ((isect_ray_tri_epsilon_v3( + ray_start, ray_normal, t0, t1, t2, &depth_test, NULL, 0.1f) && (depth_test < *depth))) + { + *depth = depth_test; + return true; + } + else { + return false; + } +} + +/* Take advantage of the fact we know this wont be an intersection. + * Just handle ray-tri edges. */ +static float dist_squared_ray_to_tri_v3_fast( + const float ray_origin[3], const float ray_direction[3], + const float v0[3], const float v1[3], const float v2[3], + float r_point[3], float *r_depth) +{ + const float *tri[3] = {v0, v1, v2}; + float dist_sq_best = FLT_MAX; + for (int i = 0, j = 2; i < 3; j = i++) { + float point_test[3], depth_test = FLT_MAX; + const float dist_sq_test = dist_squared_ray_to_seg_v3( + ray_origin, ray_direction, tri[i], tri[j], point_test, &depth_test); + if (dist_sq_test < dist_sq_best || i == 0) { + copy_v3_v3(r_point, point_test); + *r_depth = depth_test; + dist_sq_best = dist_sq_test; + } + } + return dist_sq_best; +} + +bool ray_face_nearest_quad( + const float ray_start[3], const float ray_normal[3], + const float t0[3], const float t1[3], const float t2[3], const float t3[3], + float *depth, float *dist_sq) +{ + float dist_sq_test; + float co[3], depth_test; + + if (((dist_sq_test = dist_squared_ray_to_tri_v3_fast( + ray_start, ray_normal, t0, t1, t2, co, &depth_test)) < *dist_sq)) + { + *dist_sq = dist_sq_test; + *depth = depth_test; + if (((dist_sq_test = dist_squared_ray_to_tri_v3_fast( + ray_start, ray_normal, t0, t2, t3, co, &depth_test)) < *dist_sq)) + { + *dist_sq = dist_sq_test; + *depth = depth_test; + } + return true; + } + else { + return false; + } +} + +bool ray_face_nearest_tri( + const float ray_start[3], const float ray_normal[3], + const float t0[3], const float t1[3], const float t2[3], + float *depth, float *dist_sq) { - float dist_test; + float dist_sq_test; + float co[3], depth_test; - if ((isect_ray_tri_epsilon_v3(ray_start, ray_normal, t0, t1, t2, &dist_test, NULL, 0.1f) && (dist_test < *dist))) { - *dist = dist_test; + if (((dist_sq_test = dist_squared_ray_to_tri_v3_fast( + ray_start, ray_normal, t0, t1, t2, co, &depth_test)) < *dist_sq)) + { + *dist_sq = dist_sq_test; + *depth = depth_test; return true; } else { @@ -1574,7 +1648,7 @@ static bool pbvh_faces_node_raycast( PBVH *bvh, const PBVHNode *node, float (*origco)[3], const float ray_start[3], const float ray_normal[3], - float *dist) + float *depth) { const MVert *vert = bvh->verts; const MLoop *mloop = bvh->mloop; @@ -1596,7 +1670,7 @@ static bool pbvh_faces_node_raycast( origco[face_verts[0]], origco[face_verts[1]], origco[face_verts[2]], - dist); + depth); } else { /* intersect with current coordinates */ @@ -1605,7 +1679,7 @@ static bool pbvh_faces_node_raycast( vert[mloop[lt->tri[0]].v].co, vert[mloop[lt->tri[1]].v].co, vert[mloop[lt->tri[2]].v].co, - dist); + depth); } } @@ -1616,7 +1690,7 @@ static bool pbvh_grids_node_raycast( PBVH *bvh, PBVHNode *node, float (*origco)[3], const float ray_start[3], const float ray_normal[3], - float *dist) + float *depth) { const int totgrid = node->totprim; const int gridsize = bvh->gridkey.grid_size; @@ -1646,7 +1720,7 @@ static bool pbvh_grids_node_raycast( origco[y * gridsize + x + 1], origco[(y + 1) * gridsize + x + 1], origco[(y + 1) * gridsize + x], - dist); + depth); } else { hit |= ray_face_intersection_quad( @@ -1655,7 +1729,7 @@ static bool pbvh_grids_node_raycast( CCG_grid_elem_co(&bvh->gridkey, grid, x + 1, y), CCG_grid_elem_co(&bvh->gridkey, grid, x + 1, y + 1), CCG_grid_elem_co(&bvh->gridkey, grid, x, y + 1), - dist); + depth); } } } @@ -1670,7 +1744,7 @@ static bool pbvh_grids_node_raycast( bool BKE_pbvh_node_raycast( PBVH *bvh, PBVHNode *node, float (*origco)[3], bool use_origco, const float ray_start[3], const float ray_normal[3], - float *dist) + float *depth) { bool hit = false; @@ -1681,16 +1755,16 @@ bool BKE_pbvh_node_raycast( case PBVH_FACES: hit |= pbvh_faces_node_raycast( bvh, node, origco, - ray_start, ray_normal, dist); + ray_start, ray_normal, depth); break; case PBVH_GRIDS: hit |= pbvh_grids_node_raycast( bvh, node, origco, - ray_start, ray_normal, dist); + ray_start, ray_normal, depth); break; case PBVH_BMESH: hit = pbvh_bmesh_node_raycast( - node, ray_start, ray_normal, dist, use_origco); + node, ray_start, ray_normal, depth, use_origco); break; } @@ -1741,6 +1815,176 @@ void BKE_pbvh_raycast_project_ray_root( } } +/* -------------------------------------------------------------------- */ + +typedef struct { + struct DistRayAABB_Precalc dist_ray_to_aabb_precalc; + bool original; +} FindNearestRayData; + +static bool nearest_to_ray_aabb_dist_sq(PBVHNode *node, void *data_v) +{ + FindNearestRayData *rcd = data_v; + const float *bb_min, *bb_max; + + if (rcd->original) { + /* BKE_pbvh_node_get_original_BB */ + bb_min = node->orig_vb.bmin; + bb_max = node->orig_vb.bmax; + } + else { + /* BKE_pbvh_node_get_BB */ + bb_min = node->vb.bmin; + bb_max = node->vb.bmax; + } + + float co_dummy[3], depth; + node->tmin = dist_squared_ray_to_aabb_v3(&rcd->dist_ray_to_aabb_precalc, bb_min, bb_max, co_dummy, &depth); + /* Ideally we would skip distances outside the range. */ + return depth > 0.0f; +} + +void BKE_pbvh_find_nearest_to_ray( + PBVH *bvh, BKE_pbvh_SearchNearestCallback cb, void *data, + const float ray_start[3], const float ray_normal[3], + bool original) +{ + FindNearestRayData ncd; + + dist_squared_ray_to_aabb_v3_precalc(&ncd.dist_ray_to_aabb_precalc, ray_start, ray_normal); + ncd.original = original; + + BKE_pbvh_search_callback_occluded(bvh, nearest_to_ray_aabb_dist_sq, &ncd, cb, data); +} + + +static bool pbvh_faces_node_nearest_to_ray( + PBVH *bvh, const PBVHNode *node, + float (*origco)[3], + const float ray_start[3], const float ray_normal[3], + float *depth, float *dist_sq) +{ + const MVert *vert = bvh->verts; + const MLoop *mloop = bvh->mloop; + const int *faces = node->prim_indices; + int i, totface = node->totprim; + bool hit = false; + + for (i = 0; i < totface; ++i) { + const MLoopTri *lt = &bvh->looptri[faces[i]]; + const int *face_verts = node->face_vert_indices[i]; + + if (paint_is_face_hidden(lt, vert, mloop)) + continue; + + if (origco) { + /* intersect with backuped original coordinates */ + hit |= ray_face_nearest_tri( + ray_start, ray_normal, + origco[face_verts[0]], + origco[face_verts[1]], + origco[face_verts[2]], + depth, dist_sq); + } + else { + /* intersect with current coordinates */ + hit |= ray_face_nearest_tri( + ray_start, ray_normal, + vert[mloop[lt->tri[0]].v].co, + vert[mloop[lt->tri[1]].v].co, + vert[mloop[lt->tri[2]].v].co, + depth, dist_sq); + } + } + + return hit; +} + +static bool pbvh_grids_node_nearest_to_ray( + PBVH *bvh, PBVHNode *node, + float (*origco)[3], + const float ray_start[3], const float ray_normal[3], + float *depth, float *dist_sq) +{ + const int totgrid = node->totprim; + const int gridsize = bvh->gridkey.grid_size; + bool hit = false; + + for (int i = 0; i < totgrid; ++i) { + CCGElem *grid = bvh->grids[node->prim_indices[i]]; + BLI_bitmap *gh; + + if (!grid) + continue; + + gh = bvh->grid_hidden[node->prim_indices[i]]; + + for (int y = 0; y < gridsize - 1; ++y) { + for (int x = 0; x < gridsize - 1; ++x) { + /* check if grid face is hidden */ + if (gh) { + if (paint_is_grid_face_hidden(gh, gridsize, x, y)) + continue; + } + + if (origco) { + hit |= ray_face_nearest_quad( + ray_start, ray_normal, + origco[y * gridsize + x], + origco[y * gridsize + x + 1], + origco[(y + 1) * gridsize + x + 1], + origco[(y + 1) * gridsize + x], + depth, dist_sq); + } + else { + hit |= ray_face_nearest_quad( + ray_start, ray_normal, + CCG_grid_elem_co(&bvh->gridkey, grid, x, y), + CCG_grid_elem_co(&bvh->gridkey, grid, x + 1, y), + CCG_grid_elem_co(&bvh->gridkey, grid, x + 1, y + 1), + CCG_grid_elem_co(&bvh->gridkey, grid, x, y + 1), + depth, dist_sq); + } + } + } + + if (origco) + origco += gridsize * gridsize; + } + + return hit; +} + +bool BKE_pbvh_node_find_nearest_to_ray( + PBVH *bvh, PBVHNode *node, float (*origco)[3], bool use_origco, + const float ray_start[3], const float ray_normal[3], + float *depth, float *dist_sq) +{ + bool hit = false; + + if (node->flag & PBVH_FullyHidden) + return false; + + switch (bvh->type) { + case PBVH_FACES: + hit |= pbvh_faces_node_nearest_to_ray( + bvh, node, origco, + ray_start, ray_normal, depth, dist_sq); + break; + case PBVH_GRIDS: + hit |= pbvh_grids_node_nearest_to_ray( + bvh, node, origco, + ray_start, ray_normal, depth, dist_sq); + break; + case PBVH_BMESH: + hit = pbvh_bmesh_node_nearest_to_ray( + node, ray_start, ray_normal, depth, dist_sq, use_origco); + break; + } + + return hit; +} + typedef struct { DMSetMaterial setMaterial; bool wireframe; diff --git a/source/blender/blenkernel/intern/pbvh_bmesh.c b/source/blender/blenkernel/intern/pbvh_bmesh.c index c5e49883dc6..187891e7210 100644 --- a/source/blender/blenkernel/intern/pbvh_bmesh.c +++ b/source/blender/blenkernel/intern/pbvh_bmesh.c @@ -718,20 +718,24 @@ static void pbvh_bmesh_node_drop_orig(PBVHNode *node) /****************************** EdgeQueue *****************************/ -typedef struct { +struct EdgeQueue; + +typedef struct EdgeQueue { Heap *heap; const float *center; + float center_proj[3]; /* for when we use projected coords. */ float radius_squared; float limit_len_squared; #ifdef USE_EDGEQUEUE_EVEN_SUBDIV float limit_len; #endif -#ifdef USE_EDGEQUEUE_FRONTFACE + bool (*edge_queue_tri_in_range)(const struct EdgeQueue *q, BMFace *f); + const float *view_normal; +#ifdef USE_EDGEQUEUE_FRONTFACE unsigned int use_view_normal : 1; #endif - } EdgeQueue; typedef struct { @@ -785,7 +789,6 @@ static bool edge_queue_tri_in_sphere(const EdgeQueue *q, BMFace *f) float c[3]; /* Get closest point in triangle to sphere center */ - // BM_iter_as_array(NULL, BM_VERTS_OF_FACE, f, (void **)v_tri, 3); BM_face_as_array_vert_tri(f, v_tri); closest_on_tri_to_point_v3(c, q->center, v_tri[0]->co, v_tri[1]->co, v_tri[2]->co); @@ -794,6 +797,25 @@ static bool edge_queue_tri_in_sphere(const EdgeQueue *q, BMFace *f) return len_squared_v3v3(q->center, c) <= q->radius_squared; } +static bool edge_queue_tri_in_circle(const EdgeQueue *q, BMFace *f) +{ + BMVert *v_tri[3]; + float c[3]; + float tri_proj[3][3]; + + /* Get closest point in triangle to sphere center */ + BM_face_as_array_vert_tri(f, v_tri); + + project_plane_normalized_v3_v3v3(tri_proj[0], v_tri[0]->co, q->view_normal); + project_plane_normalized_v3_v3v3(tri_proj[1], v_tri[1]->co, q->view_normal); + project_plane_normalized_v3_v3v3(tri_proj[2], v_tri[2]->co, q->view_normal); + + closest_on_tri_to_point_v3(c, q->center_proj, tri_proj[0], tri_proj[1], tri_proj[2]); + + /* Check if triangle intersects the sphere */ + return len_squared_v3v3(q->center_proj, c) <= q->radius_squared; +} + /* Return true if the vertex mask is less than 1.0, false otherwise */ static bool check_mask(EdgeQueueContext *eq_ctx, BMVert *v) { @@ -929,7 +951,7 @@ static void long_edge_queue_face_add( } #endif - if (edge_queue_tri_in_sphere(eq_ctx->q, f)) { + if (eq_ctx->q->edge_queue_tri_in_range(eq_ctx->q, f)) { /* Check each edge of the face */ BMLoop *l_first = BM_FACE_FIRST_LOOP(f); BMLoop *l_iter = l_first; @@ -960,7 +982,7 @@ static void short_edge_queue_face_add( } #endif - if (edge_queue_tri_in_sphere(eq_ctx->q, f)) { + if (eq_ctx->q->edge_queue_tri_in_range(eq_ctx->q, f)) { BMLoop *l_iter; BMLoop *l_first; @@ -984,7 +1006,7 @@ static void short_edge_queue_face_add( static void long_edge_queue_create( EdgeQueueContext *eq_ctx, PBVH *bvh, const float center[3], const float view_normal[3], - float radius) + float radius, const bool use_frontface, const bool use_projected) { eq_ctx->q->heap = BLI_heap_new(); eq_ctx->q->center = center; @@ -994,13 +1016,22 @@ static void long_edge_queue_create( eq_ctx->q->limit_len = bvh->bm_max_edge_len; #endif -#ifdef USE_EDGEQUEUE_FRONTFACE eq_ctx->q->view_normal = view_normal; - eq_ctx->q->use_view_normal = (view_normal != NULL); + +#ifdef USE_EDGEQUEUE_FRONTFACE + eq_ctx->q->use_view_normal = use_frontface; #else - UNUSED_VARS(view_normal); + UNUSED_VARS(use_frontface); #endif + if (use_projected) { + eq_ctx->q->edge_queue_tri_in_range = edge_queue_tri_in_circle; + project_plane_normalized_v3_v3v3(eq_ctx->q->center_proj, center, view_normal); + } + else { + eq_ctx->q->edge_queue_tri_in_range = edge_queue_tri_in_sphere; + } + #ifdef USE_EDGEQUEUE_TAG_VERIFY pbvh_bmesh_edge_tag_verify(bvh); #endif @@ -1037,7 +1068,7 @@ static void long_edge_queue_create( static void short_edge_queue_create( EdgeQueueContext *eq_ctx, PBVH *bvh, const float center[3], const float view_normal[3], - float radius) + float radius, const bool use_frontface, const bool use_projected) { eq_ctx->q->heap = BLI_heap_new(); eq_ctx->q->center = center; @@ -1047,13 +1078,22 @@ static void short_edge_queue_create( eq_ctx->q->limit_len = bvh->bm_min_edge_len; #endif -#ifdef USE_EDGEQUEUE_FRONTFACE eq_ctx->q->view_normal = view_normal; - eq_ctx->q->use_view_normal = (view_normal != NULL); + +#ifdef USE_EDGEQUEUE_FRONTFACE + eq_ctx->q->use_view_normal = use_frontface; #else - UNUSED_VARS(view_normal); + UNUSED_VARS(use_frontface); #endif + if (use_projected) { + eq_ctx->q->edge_queue_tri_in_range = edge_queue_tri_in_circle; + project_plane_normalized_v3_v3v3(eq_ctx->q->center_proj, center, view_normal); + } + else { + eq_ctx->q->edge_queue_tri_in_range = edge_queue_tri_in_sphere; + } + for (int n = 0; n < bvh->totnode; n++) { PBVHNode *node = &bvh->nodes[n]; @@ -1466,7 +1506,7 @@ static bool pbvh_bmesh_collapse_short_edges( bool pbvh_bmesh_node_raycast( PBVHNode *node, const float ray_start[3], - const float ray_normal[3], float *dist, + const float ray_normal[3], float *depth, bool use_original) { bool hit = false; @@ -1479,7 +1519,7 @@ bool pbvh_bmesh_node_raycast( node->bm_orco[t[0]], node->bm_orco[t[1]], node->bm_orco[t[2]], - dist); + depth); } } else { @@ -1498,7 +1538,7 @@ bool pbvh_bmesh_node_raycast( v_tri[0]->co, v_tri[1]->co, v_tri[2]->co, - dist); + depth); } } } @@ -1509,7 +1549,7 @@ bool pbvh_bmesh_node_raycast( bool BKE_pbvh_bmesh_node_raycast_detail( PBVHNode *node, const float ray_start[3], const float ray_normal[3], - float *dist, float *r_detail) + float *depth, float *r_detail) { if (node->flag & PBVH_FullyHidden) return 0; @@ -1531,7 +1571,7 @@ bool BKE_pbvh_bmesh_node_raycast_detail( v_tri[0]->co, v_tri[1]->co, v_tri[2]->co, - dist); + depth); if (hit_local) { f_hit = f; @@ -1554,6 +1594,47 @@ bool BKE_pbvh_bmesh_node_raycast_detail( return hit; } +bool pbvh_bmesh_node_nearest_to_ray( + PBVHNode *node, const float ray_start[3], + const float ray_normal[3], float *depth, float *dist_sq, + bool use_original) +{ + bool hit = false; + + if (use_original && node->bm_tot_ortri) { + for (int i = 0; i < node->bm_tot_ortri; i++) { + const int *t = node->bm_ortri[i]; + hit |= ray_face_nearest_tri( + ray_start, ray_normal, + node->bm_orco[t[0]], + node->bm_orco[t[1]], + node->bm_orco[t[2]], + depth, dist_sq); + } + } + else { + GSetIterator gs_iter; + + GSET_ITER (gs_iter, node->bm_faces) { + BMFace *f = BLI_gsetIterator_getKey(&gs_iter); + + BLI_assert(f->len == 3); + if (!BM_elem_flag_test(f, BM_ELEM_HIDDEN)) { + BMVert *v_tri[3]; + + BM_face_as_array_vert_tri(f, v_tri); + hit |= ray_face_nearest_tri( + ray_start, ray_normal, + v_tri[0]->co, + v_tri[1]->co, + v_tri[2]->co, + depth, dist_sq); + } + } + } + + return hit; +} void pbvh_bmesh_normals_update(PBVHNode **nodes, int totnode) { @@ -1854,7 +1935,7 @@ void BKE_pbvh_build_bmesh( bool BKE_pbvh_bmesh_update_topology( PBVH *bvh, PBVHTopologyUpdateMode mode, const float center[3], const float view_normal[3], - float radius) + float radius, const bool use_frontface, const bool use_projected) { /* 2 is enough for edge faces - manifold edge */ BLI_buffer_declare_static(BMLoop *, edge_loops, BLI_BUFFER_NOP, 2); @@ -1877,7 +1958,7 @@ bool BKE_pbvh_bmesh_update_topology( cd_vert_mask_offset, cd_vert_node_offset, cd_face_node_offset, }; - short_edge_queue_create(&eq_ctx, bvh, center, view_normal, radius); + short_edge_queue_create(&eq_ctx, bvh, center, view_normal, radius, use_frontface, use_projected); modified |= pbvh_bmesh_collapse_short_edges( &eq_ctx, bvh, &deleted_faces); BLI_heap_free(q.heap, NULL); @@ -1892,7 +1973,7 @@ bool BKE_pbvh_bmesh_update_topology( cd_vert_mask_offset, cd_vert_node_offset, cd_face_node_offset, }; - long_edge_queue_create(&eq_ctx, bvh, center, view_normal, radius); + long_edge_queue_create(&eq_ctx, bvh, center, view_normal, radius, use_frontface, use_projected); modified |= pbvh_bmesh_subdivide_long_edges( &eq_ctx, bvh, &edge_loops); BLI_heap_free(q.heap, NULL); diff --git a/source/blender/blenkernel/intern/pbvh_intern.h b/source/blender/blenkernel/intern/pbvh_intern.h index 01057318568..e05a3068682 100644 --- a/source/blender/blenkernel/intern/pbvh_intern.h +++ b/source/blender/blenkernel/intern/pbvh_intern.h @@ -185,11 +185,21 @@ void pbvh_grow_nodes(PBVH *bvh, int totnode); bool ray_face_intersection_quad( const float ray_start[3], const float ray_normal[3], const float *t0, const float *t1, const float *t2, const float *t3, - float *r_dist); + float *depth); bool ray_face_intersection_tri( const float ray_start[3], const float ray_normal[3], const float *t0, const float *t1, const float *t2, - float *r_dist); + float *depth); + +bool ray_face_nearest_quad( + const float ray_start[3], const float ray_normal[3], + const float *t0, const float *t1, const float *t2, const float *t3, + float *r_depth, float *r_dist_sq); +bool ray_face_nearest_tri( + const float ray_start[3], const float ray_normal[3], + const float *t0, const float *t1, const float *t2, + float *r_depth, float *r_dist_sq); + void pbvh_update_BB_redraw(PBVH *bvh, PBVHNode **nodes, int totnode, int flag); /* pbvh_bmesh.c */ @@ -197,6 +207,10 @@ bool pbvh_bmesh_node_raycast( PBVHNode *node, const float ray_start[3], const float ray_normal[3], float *dist, bool use_original); +bool pbvh_bmesh_node_nearest_to_ray( + PBVHNode *node, const float ray_start[3], + const float ray_normal[3], float *depth, float *dist_sq, + bool use_original); void pbvh_bmesh_normals_update(PBVHNode **nodes, int totnode); diff --git a/source/blender/blenkernel/intern/scene.c b/source/blender/blenkernel/intern/scene.c index 1b0e09afcb5..fca0ff92a45 100644 --- a/source/blender/blenkernel/intern/scene.c +++ b/source/blender/blenkernel/intern/scene.c @@ -2019,7 +2019,7 @@ int BKE_scene_num_threads(const Scene *scene) int BKE_render_preview_pixel_size(const RenderData *r) { if (r->preview_pixel_size == 0) { - return (U.pixelsize > 1.5f)? 2 : 1; + return (U.pixelsize > 1.5f) ? 2 : 1; } return r->preview_pixel_size; } diff --git a/source/blender/blenkernel/intern/seqeffects.c b/source/blender/blenkernel/intern/seqeffects.c index e435d87024e..a2c45057bf7 100644 --- a/source/blender/blenkernel/intern/seqeffects.c +++ b/source/blender/blenkernel/intern/seqeffects.c @@ -1861,7 +1861,7 @@ static void RVBlurBitmap2_float(float *map, int width, int height, float blur, i float *filter = NULL; int x, y, i, fx, fy; int index, ix, halfWidth; - float fval, k, curColor[3], curColor2[3], weight = 0; + float fval, k, curColor[4], curColor2[4], weight = 0; /* If we're not really blurring, bail out */ if (blur <= 0) @@ -1906,47 +1906,38 @@ static void RVBlurBitmap2_float(float *map, int width, int height, float blur, i for (y = 0; y < height; y++) { /* Do the left & right strips */ for (x = 0; x < halfWidth; x++) { - index = (x + y * width) * 4; fx = 0; - curColor[0] = curColor[1] = curColor[2] = 0.0f; - curColor2[0] = curColor2[1] = curColor2[2] = 0.0f; + zero_v4(curColor); + zero_v4(curColor2); for (i = x - halfWidth; i < x + halfWidth; i++) { if ((i >= 0) && (i < width)) { - curColor[0] += map[(i + y * width) * 4 + GlowR] * filter[fx]; - curColor[1] += map[(i + y * width) * 4 + GlowG] * filter[fx]; - curColor[2] += map[(i + y * width) * 4 + GlowB] * filter[fx]; + index = (i + y * width) * 4; + madd_v4_v4fl(curColor, map + index, filter[fx]); - curColor2[0] += map[(width - 1 - i + y * width) * 4 + GlowR] * filter[fx]; - curColor2[1] += map[(width - 1 - i + y * width) * 4 + GlowG] * filter[fx]; - curColor2[2] += map[(width - 1 - i + y * width) * 4 + GlowB] * filter[fx]; + index = (width - 1 - i + y * width) * 4; + madd_v4_v4fl(curColor2, map + index, filter[fx]); } fx++; } - temp[index + GlowR] = curColor[0]; - temp[index + GlowG] = curColor[1]; - temp[index + GlowB] = curColor[2]; - - temp[((width - 1 - x + y * width) * 4) + GlowR] = curColor2[0]; - temp[((width - 1 - x + y * width) * 4) + GlowG] = curColor2[1]; - temp[((width - 1 - x + y * width) * 4) + GlowB] = curColor2[2]; + index = (x + y * width) * 4; + copy_v4_v4(temp + index, curColor); + index = (width - 1 - x + y * width) * 4; + copy_v4_v4(temp + index, curColor2); } /* Do the main body */ for (x = halfWidth; x < width - halfWidth; x++) { - index = (x + y * width) * 4; fx = 0; - zero_v3(curColor); + zero_v4(curColor); for (i = x - halfWidth; i < x + halfWidth; i++) { - curColor[0] += map[(i + y * width) * 4 + GlowR] * filter[fx]; - curColor[1] += map[(i + y * width) * 4 + GlowG] * filter[fx]; - curColor[2] += map[(i + y * width) * 4 + GlowB] * filter[fx]; + index = (i + y * width) * 4; + madd_v4_v4fl(curColor, map + index, filter[fx]); fx++; } - temp[index + GlowR] = curColor[0]; - temp[index + GlowG] = curColor[1]; - temp[index + GlowB] = curColor[2]; + index = (x + y * width) * 4; + copy_v4_v4(temp + index, curColor); } } @@ -1957,46 +1948,39 @@ static void RVBlurBitmap2_float(float *map, int width, int height, float blur, i for (x = 0; x < width; x++) { /* Do the top & bottom strips */ for (y = 0; y < halfWidth; y++) { - index = (x + y * width) * 4; fy = 0; - zero_v3(curColor); - zero_v3(curColor2); + zero_v4(curColor); + zero_v4(curColor2); for (i = y - halfWidth; i < y + halfWidth; i++) { if ((i >= 0) && (i < height)) { /* Bottom */ - curColor[0] += map[(x + i * width) * 4 + GlowR] * filter[fy]; - curColor[1] += map[(x + i * width) * 4 + GlowG] * filter[fy]; - curColor[2] += map[(x + i * width) * 4 + GlowB] * filter[fy]; + index = (x + i * width) * 4; + madd_v4_v4fl(curColor, map + index, filter[fy]); /* Top */ - curColor2[0] += map[(x + (height - 1 - i) * width) * 4 + GlowR] * filter[fy]; - curColor2[1] += map[(x + (height - 1 - i) * width) * 4 + GlowG] * filter[fy]; - curColor2[2] += map[(x + (height - 1 - i) * width) * 4 + GlowB] * filter[fy]; + index = (x + (height - 1 - i) * width) * 4; + madd_v4_v4fl(curColor2, map + index, filter[fy]); } fy++; } - temp[index + GlowR] = curColor[0]; - temp[index + GlowG] = curColor[1]; - temp[index + GlowB] = curColor[2]; - temp[((x + (height - 1 - y) * width) * 4) + GlowR] = curColor2[0]; - temp[((x + (height - 1 - y) * width) * 4) + GlowG] = curColor2[1]; - temp[((x + (height - 1 - y) * width) * 4) + GlowB] = curColor2[2]; + index = (x + y * width) * 4; + copy_v4_v4(temp + index, curColor); + + index = (x + (height - 1 - y) * width) * 4; + copy_v4_v4(temp + index, curColor2); } /* Do the main body */ for (y = halfWidth; y < height - halfWidth; y++) { - index = (x + y * width) * 4; fy = 0; - zero_v3(curColor); + zero_v4(curColor); for (i = y - halfWidth; i < y + halfWidth; i++) { - curColor[0] += map[(x + i * width) * 4 + GlowR] * filter[fy]; - curColor[1] += map[(x + i * width) * 4 + GlowG] * filter[fy]; - curColor[2] += map[(x + i * width) * 4 + GlowB] * filter[fy]; + index = (x + i * width) * 4; + madd_v4_v4fl(curColor, map + index, filter[fy]); fy++; } - temp[index + GlowR] = curColor[0]; - temp[index + GlowG] = curColor[1]; - temp[index + GlowB] = curColor[2]; + index = (x + y * width) * 4; + copy_v4_v4(temp + index, curColor); } } @@ -2015,10 +1999,10 @@ static void RVAddBitmaps_float(float *a, float *b, float *c, int width, int heig for (y = 0; y < height; y++) { for (x = 0; x < width; x++) { index = (x + y * width) * 4; - c[index + GlowR] = MIN2(1.0f, a[index + GlowR] + b[index + GlowR]); - c[index + GlowG] = MIN2(1.0f, a[index + GlowG] + b[index + GlowG]); - c[index + GlowB] = MIN2(1.0f, a[index + GlowB] + b[index + GlowB]); - c[index + GlowA] = MIN2(1.0f, a[index + GlowA] + b[index + GlowA]); + c[index + GlowR] = min_ff(1.0f, a[index + GlowR] + b[index + GlowR]); + c[index + GlowG] = min_ff(1.0f, a[index + GlowG] + b[index + GlowG]); + c[index + GlowB] = min_ff(1.0f, a[index + GlowB] + b[index + GlowB]); + c[index + GlowA] = min_ff(1.0f, a[index + GlowA] + b[index + GlowA]); } } } @@ -2035,10 +2019,10 @@ static void RVIsolateHighlights_float(float *in, float *out, int width, int heig /* Isolate the intensity */ intensity = (in[index + GlowR] + in[index + GlowG] + in[index + GlowB] - threshold); if (intensity > 0) { - out[index + GlowR] = MIN2(clamp, (in[index + GlowR] * boost * intensity)); - out[index + GlowG] = MIN2(clamp, (in[index + GlowG] * boost * intensity)); - out[index + GlowB] = MIN2(clamp, (in[index + GlowB] * boost * intensity)); - out[index + GlowA] = MIN2(clamp, (in[index + GlowA] * boost * intensity)); + out[index + GlowR] = min_ff(clamp, (in[index + GlowR] * boost * intensity)); + out[index + GlowG] = min_ff(clamp, (in[index + GlowG] * boost * intensity)); + out[index + GlowB] = min_ff(clamp, (in[index + GlowB] * boost * intensity)); + out[index + GlowA] = min_ff(clamp, (in[index + GlowA] * boost * intensity)); } else { out[index + GlowR] = 0; diff --git a/source/blender/blenkernel/intern/softbody.c b/source/blender/blenkernel/intern/softbody.c index 7cea1e4940f..1ba8fe83971 100644 --- a/source/blender/blenkernel/intern/softbody.c +++ b/source/blender/blenkernel/intern/softbody.c @@ -3414,7 +3414,7 @@ static void softbody_update_positions(Object *ob, SoftBody *sb, float (*vertexCo * lloc, lrot, lscale are allowed to be NULL, just in case you don't need it. * should be pretty useful for pythoneers :) * not! velocity .. 2nd order stuff - * vcloud_estimate_transform see + * vcloud_estimate_transform_v3 see */ void SB_estimate_transform(Object *ob, float lloc[3], float lrot[3][3], float lscale[3][3]) @@ -3438,7 +3438,7 @@ void SB_estimate_transform(Object *ob, float lloc[3], float lrot[3][3], float ls copy_v3_v3(opos[a], bp->pos); } - vcloud_estimate_transform(sb->totpoint, opos, NULL, rpos, NULL, com, rcom, lrot, lscale); + vcloud_estimate_transform_v3(sb->totpoint, opos, NULL, rpos, NULL, com, rcom, lrot, lscale); //sub_v3_v3(com, rcom); if (lloc) copy_v3_v3(lloc, com); copy_v3_v3(sb->lcom, com); diff --git a/source/blender/blenlib/BLI_math_geom.h b/source/blender/blenlib/BLI_math_geom.h index d0b59244384..933e31ba84b 100644 --- a/source/blender/blenlib/BLI_math_geom.h +++ b/source/blender/blenlib/BLI_math_geom.h @@ -126,15 +126,15 @@ struct DistRayAABB_Precalc { float ray_inv_dir[3]; bool sign[3]; }; -void dist_squared_ray_to_aabb_precalc( +void dist_squared_ray_to_aabb_v3_precalc( struct DistRayAABB_Precalc *neasrest_precalc, const float ray_origin[3], const float ray_direction[3]); -float dist_squared_ray_to_aabb( +float dist_squared_ray_to_aabb_v3( const struct DistRayAABB_Precalc *data, const float bb_min[3], const float bb_max[3], float r_point[3], float *r_depth); /* when there is no advantage to precalc. */ -float dist_squared_to_ray_to_aabb_simple( +float dist_squared_ray_to_aabb_v3_simple( const float ray_origin[3], const float ray_direction[3], const float bb_min[3], const float bb_max[3], float r_point[3], float *r_depth); @@ -407,23 +407,23 @@ void map_to_plane_axis_angle_v2_v3v3fl(float r_co[2], const float co[3], const f /********************************** Normals **********************************/ -void accumulate_vertex_normals_tri( +void accumulate_vertex_normals_tri_v3( float n1[3], float n2[3], float n3[3], const float f_no[3], const float co1[3], const float co2[3], const float co3[3]); -void accumulate_vertex_normals( +void accumulate_vertex_normals_v3( float n1[3], float n2[3], float n3[3], float n4[3], const float f_no[3], const float co1[3], const float co2[3], const float co3[3], const float co4[3]); -void accumulate_vertex_normals_poly( +void accumulate_vertex_normals_poly_v3( float **vertnos, const float polyno[3], const float **vertcos, float vdiffs[][3], const int nverts); /********************************* Tangents **********************************/ -void tangent_from_uv( +void tangent_from_uv_v3( const float uv1[2], const float uv2[2], const float uv3[2], const float co1[3], const float co2[3], const float co3[3], const float n[3], @@ -431,9 +431,9 @@ void tangent_from_uv( /******************************** Vector Clouds ******************************/ -void vcloud_estimate_transform(int list_size, float (*pos)[3], float *weight, - float (*rpos)[3], float *rweight, - float lloc[3], float rloc[3], float lrot[3][3], float lscale[3][3]); +void vcloud_estimate_transform_v3( + const int list_size, const float (*pos)[3], const float *weight, const float (*rpos)[3], const float *rweight, + float lloc[3], float rloc[3], float lrot[3][3], float lscale[3][3]); /****************************** Spherical Harmonics *************************/ @@ -464,7 +464,7 @@ float form_factor_hemi_poly(float p[3], float n[3], float v1[3], float v2[3], float v3[3], float v4[3]); void axis_dominant_v3_to_m3_negate(float r_mat[3][3], const float normal[3]); -void axis_dominant_v3_to_m3(float r_mat[3][3], const float normal[3]); +void axis_dominant_v3_to_m3(float r_mat[3][3], const float normal[3]); MINLINE void axis_dominant_v3(int *r_axis_a, int *r_axis_b, const float axis[3]); MINLINE float axis_dominant_v3_max(int *r_axis_a, int *r_axis_b, const float axis[3]) ATTR_WARN_UNUSED_RESULT; diff --git a/source/blender/blenlib/intern/math_geom.c b/source/blender/blenlib/intern/math_geom.c index dbbc1adb534..d3080e5530f 100644 --- a/source/blender/blenlib/intern/math_geom.c +++ b/source/blender/blenlib/intern/math_geom.c @@ -623,7 +623,7 @@ float dist_squared_ray_to_seg_v3( /** \name dist_squared_to_ray_to_aabb and helpers * \{ */ -void dist_squared_ray_to_aabb_precalc( +void dist_squared_ray_to_aabb_v3_precalc( struct DistRayAABB_Precalc *neasrest_precalc, const float ray_origin[3], const float ray_direction[3]) { @@ -641,7 +641,7 @@ void dist_squared_ray_to_aabb_precalc( /** * Returns the distance from a ray to a bound-box (projected on ray) */ -float dist_squared_ray_to_aabb( +float dist_squared_ray_to_aabb_v3( const struct DistRayAABB_Precalc *data, const float bb_min[3], const float bb_max[3], float r_point[3], float *r_depth) @@ -753,14 +753,14 @@ float dist_squared_ray_to_aabb( r_point, r_depth); } -float dist_squared_to_ray_to_aabb_simple( +float dist_squared_ray_to_aabb_v3_simple( const float ray_origin[3], const float ray_direction[3], const float bbmin[3], const float bbmax[3], float r_point[3], float *r_depth) { struct DistRayAABB_Precalc data; - dist_squared_ray_to_aabb_precalc(&data, ray_origin, ray_direction); - return dist_squared_ray_to_aabb(&data, bbmin, bbmax, r_point, r_depth); + dist_squared_ray_to_aabb_v3_precalc(&data, ray_origin, ray_direction); + return dist_squared_ray_to_aabb_v3(&data, bbmin, bbmax, r_point, r_depth); } /** \} */ @@ -4053,7 +4053,7 @@ void map_to_plane_axis_angle_v2_v3v3fl(float r_co[2], const float co[3], const f /********************************* Normals **********************************/ -void accumulate_vertex_normals_tri( +void accumulate_vertex_normals_tri_v3( float n1[3], float n2[3], float n3[3], const float f_no[3], const float co1[3], const float co2[3], const float co3[3]) @@ -4087,7 +4087,7 @@ void accumulate_vertex_normals_tri( } } -void accumulate_vertex_normals( +void accumulate_vertex_normals_v3( float n1[3], float n2[3], float n3[3], float n4[3], const float f_no[3], const float co1[3], const float co2[3], const float co3[3], const float co4[3]) @@ -4131,7 +4131,7 @@ void accumulate_vertex_normals( /* Add weighted face normal component into normals of the face vertices. * Caller must pass pre-allocated vdiffs of nverts length. */ -void accumulate_vertex_normals_poly(float **vertnos, const float polyno[3], +void accumulate_vertex_normals_poly_v3(float **vertnos, const float polyno[3], const float **vertcos, float vdiffs[][3], const int nverts) { int i; @@ -4162,7 +4162,7 @@ void accumulate_vertex_normals_poly(float **vertnos, const float polyno[3], /********************************* Tangents **********************************/ -void tangent_from_uv( +void tangent_from_uv_v3( const float uv1[2], const float uv2[2], const float uv3[3], const float co1[3], const float co2[3], const float co3[3], const float n[3], @@ -4204,30 +4204,28 @@ void tangent_from_uv( /****************************** Vector Clouds ********************************/ /* vector clouds */ -/* void vcloud_estimate_transform(int list_size, float (*pos)[3], float *weight, float (*rpos)[3], float *rweight, - * float lloc[3], float rloc[3], float lrot[3][3], float lscale[3][3]) - * +/** * input - * ( - * int list_size - * 4 lists as pointer to array[list_size] - * 1. current pos array of 'new' positions - * 2. current weight array of 'new'weights (may be NULL pointer if you have no weights ) - * 3. reference rpos array of 'old' positions - * 4. reference rweight array of 'old'weights (may be NULL pointer if you have no weights ) - * ) + * + * \param list_size: 4 lists as pointer to array[list_size] + * \param pos: current pos array of 'new' positions + * \param weight: current weight array of 'new'weights (may be NULL pointer if you have no weights) + * \param rpos: Reference rpos array of 'old' positions + * \param rweight: Reference rweight array of 'old'weights (may be NULL pointer if you have no weights). + * * output - * ( - * float lloc[3] center of mass pos - * float rloc[3] center of mass rpos - * float lrot[3][3] rotation matrix - * float lscale[3][3] scale matrix + * + * \param lloc: Center of mass pos. + * \param rloc: Center of mass rpos. + * \param lrot: Rotation matrix. + * \param lscale: Scale matrix. + * * pointers may be NULL if not needed - * ) */ -void vcloud_estimate_transform(int list_size, float (*pos)[3], float *weight, float (*rpos)[3], float *rweight, - float lloc[3], float rloc[3], float lrot[3][3], float lscale[3][3]) +void vcloud_estimate_transform_v3( + const int list_size, const float (*pos)[3], const float *weight, const float (*rpos)[3], const float *rweight, + float lloc[3], float rloc[3], float lrot[3][3], float lscale[3][3]) { float accu_com[3] = {0.0f, 0.0f, 0.0f}, accu_rcom[3] = {0.0f, 0.0f, 0.0f}; float accu_weight = 0.0f, accu_rweight = 0.0f; diff --git a/source/blender/blenlib/intern/noise.c b/source/blender/blenlib/intern/noise.c index 86c24307ae2..83012694ac0 100644 --- a/source/blender/blenlib/intern/noise.c +++ b/source/blender/blenlib/intern/noise.c @@ -1395,9 +1395,9 @@ static float voronoi_CrS(float x, float y, float z) static float cellNoiseU(float x, float y, float z) { /* avoid precision issues on unit coordinates */ - x = (x + 0.000001f)*1.00001f; - y = (y + 0.000001f)*1.00001f; - z = (z + 0.000001f)*1.00001f; + x = (x + 0.000001f) * 1.00001f; + y = (y + 0.000001f) * 1.00001f; + z = (z + 0.000001f) * 1.00001f; int xi = (int)(floor(x)); int yi = (int)(floor(y)); @@ -1417,9 +1417,9 @@ float cellNoise(float x, float y, float z) void cellNoiseV(float x, float y, float z, float ca[3]) { /* avoid precision issues on unit coordinates */ - x = (x + 0.000001f)*1.00001f; - y = (y + 0.000001f)*1.00001f; - z = (z + 0.000001f)*1.00001f; + x = (x + 0.000001f) * 1.00001f; + y = (y + 0.000001f) * 1.00001f; + z = (z + 0.000001f) * 1.00001f; int xi = (int)(floor(x)); int yi = (int)(floor(y)); diff --git a/source/blender/blenloader/intern/versioning_270.c b/source/blender/blenloader/intern/versioning_270.c index 86c5a7913d3..b945f5bdcd3 100644 --- a/source/blender/blenloader/intern/versioning_270.c +++ b/source/blender/blenloader/intern/versioning_270.c @@ -1698,13 +1698,21 @@ void blo_do_versions_270(FileData *fd, Library *UNUSED(lib), Main *main) /* Fix for invalid state of screen due to bug in older versions. */ for (bScreen *sc = main->screen.first; sc; sc = sc->id.next) { for (ScrArea *sa = sc->areabase.first; sa; sa = sa->next) { - if(sa->full && sc->state == SCREENNORMAL) { + if (sa->full && sc->state == SCREENNORMAL) { sa->full = NULL; } } } - if (!DNA_struct_elem_find(fd->filesdna, "VPaint", "char", "falloff_shape")) { + if (!DNA_struct_elem_find(fd->filesdna, "Brush", "float", "falloff_angle")) { + for (Brush *br = main->brush.first; br; br = br->id.next) { + br->falloff_angle = DEG2RADF(80); + br->flag &= ~( + BRUSH_FLAG_DEPRECATED_1 | BRUSH_FLAG_DEPRECATED_2 | + BRUSH_FLAG_DEPRECATED_3 | BRUSH_FLAG_DEPRECATED_4 | + BRUSH_FRONTFACE_FALLOFF); + } + for (Scene *scene = main->scene.first; scene; scene = scene->id.next) { ToolSettings *ts = scene->toolsettings; for (int i = 0; i < 2; i++) { @@ -1712,7 +1720,6 @@ void blo_do_versions_270(FileData *fd, Library *UNUSED(lib), Main *main) if (vp != NULL) { /* remove all other flags */ vp->flag &= (VP_FLAG_VGROUP_RESTRICT); - vp->normal_angle = 80; } } } diff --git a/source/blender/bmesh/intern/bmesh_mesh.c b/source/blender/bmesh/intern/bmesh_mesh.c index d5d9e4abe2c..2ff670c770e 100644 --- a/source/blender/bmesh/intern/bmesh_mesh.c +++ b/source/blender/bmesh/intern/bmesh_mesh.c @@ -775,7 +775,7 @@ static void bm_mesh_loops_calc_normals( } { - /* Code similar to accumulate_vertex_normals_poly. */ + /* Code similar to accumulate_vertex_normals_poly_v3. */ /* Calculate angle between the two poly edges incident on this vertex. */ const BMFace *f = lfan_pivot->f; const float fac = saacos(dot_v3v3(vec_next, vec_curr)); diff --git a/source/blender/editors/interface/interface_handlers.c b/source/blender/editors/interface/interface_handlers.c index b11ff27f446..9b3e98c2ace 100644 --- a/source/blender/editors/interface/interface_handlers.c +++ b/source/blender/editors/interface/interface_handlers.c @@ -9245,11 +9245,17 @@ static int ui_handle_menu_event( doit = true; } } - else if (count == act) { + else if (ELEM(but->type, + UI_BTYPE_BUT, + UI_BTYPE_BUT_MENU, + UI_BTYPE_MENU, UI_BTYPE_BLOCK, + UI_BTYPE_PULLDOWN) && + count == act) + { doit = true; } - if (doit) { + if (!(but->flag & UI_BUT_DISABLED) && doit) { /* activate buttons but open menu's */ uiButtonActivateType activate; if (but->type == UI_BTYPE_PULLDOWN) { @@ -9303,8 +9309,7 @@ static int ui_handle_menu_event( break; for (but = block->buttons.first; but; but = but->next) { - - if (but->menu_key == event->type) { + if (!(but->flag & UI_BUT_DISABLED) && but->menu_key == event->type) { if (ELEM(but->type, UI_BTYPE_BUT, UI_BTYPE_BUT_MENU)) { /* mainly for operator buttons */ ui_handle_button_activate(C, ar, but, BUTTON_ACTIVATE_APPLY); @@ -10225,9 +10230,9 @@ void UI_popup_handlers_remove(ListBase *handlers, uiPopupBlockHandle *popup) handler->ui_userdata == popup) { /* tag refresh parent popup */ - if (handler->next && - handler->next->ui_handle == ui_popup_handler && - handler->next->ui_remove == ui_popup_handler_remove) + if (handler->next && + handler->next->ui_handle == ui_popup_handler && + handler->next->ui_remove == ui_popup_handler_remove) { uiPopupBlockHandle *parent_popup = handler->next->ui_userdata; ED_region_tag_refresh_ui(parent_popup->region); diff --git a/source/blender/editors/render/render_internal.c b/source/blender/editors/render/render_internal.c index 212d6128a9e..ababd637e00 100644 --- a/source/blender/editors/render/render_internal.c +++ b/source/blender/editors/render/render_internal.c @@ -1321,7 +1321,7 @@ static void render_view3d_startjob(void *customdata, short *stop, short *do_upda if (restore) RE_DataBase_IncrementalView(re, rp->viewmat, 1); - rp->resolution_divider = MAX2(rp->resolution_divider/2, pixel_size); + rp->resolution_divider = MAX2(rp->resolution_divider / 2, pixel_size); *do_update = 1; render_update_resolution(re, rp, use_border, &cliprct); diff --git a/source/blender/editors/screen/screen_context.c b/source/blender/editors/screen/screen_context.c index 5f9a54a4654..5813f317295 100644 --- a/source/blender/editors/screen/screen_context.c +++ b/source/blender/editors/screen/screen_context.c @@ -610,8 +610,7 @@ int ed_screen_context(const bContext *C, const char *member, bContextDataResult return 1; } } - else if (CTX_data_equals(member, "selected_editable_fcurves")) - { + else if (CTX_data_equals(member, "selected_editable_fcurves")) { bAnimContext ac; if (ANIM_animdata_get_context(C, &ac) && ELEM(ac.spacetype, SPACE_ACTION, SPACE_IPO)) { diff --git a/source/blender/editors/sculpt_paint/paint_cursor.c b/source/blender/editors/sculpt_paint/paint_cursor.c index 04c7e6e2a62..befb635cfc4 100644 --- a/source/blender/editors/sculpt_paint/paint_cursor.c +++ b/source/blender/editors/sculpt_paint/paint_cursor.c @@ -42,6 +42,7 @@ #include "DNA_scene_types.h" #include "DNA_screen_types.h" #include "DNA_userdef_types.h" +#include "DNA_view3d_types.h" #include "BKE_brush.h" #include "BKE_context.h" @@ -1024,9 +1025,9 @@ static void paint_draw_cursor(bContext *C, int x, int y, void *UNUSED(unused)) ViewContext vc; view3d_set_viewcontext(C, &vc); - float zoomx, zoomy; - get_imapaint_zoom(C, &zoomx, &zoomy); - zoomx = max_ff(zoomx, zoomy); + if (vc.rv3d->rflag & RV3D_NAVIGATING) { + return; + } /* skip everything and draw brush here */ if (brush->flag & BRUSH_CURVE) { @@ -1034,6 +1035,10 @@ static void paint_draw_cursor(bContext *C, int x, int y, void *UNUSED(unused)) return; } + float zoomx, zoomy; + get_imapaint_zoom(C, &zoomx, &zoomy); + zoomx = max_ff(zoomx, zoomy); + /* set various defaults */ const float *outline_col = brush->add_col; const float outline_alpha = 0.5f; @@ -1064,11 +1069,8 @@ static void paint_draw_cursor(bContext *C, int x, int y, void *UNUSED(unused)) /* check if brush is subtracting, use different color then */ /* TODO: no way currently to know state of pen flip or * invert key modifier without starting a stroke */ - if (((ups->draw_inverted == 0) ^ - ((brush->flag & BRUSH_DIR_IN) == 0)) && - ELEM(brush->sculpt_tool, SCULPT_TOOL_DRAW, - SCULPT_TOOL_INFLATE, SCULPT_TOOL_CLAY, - SCULPT_TOOL_PINCH, SCULPT_TOOL_CREASE)) + if (((ups->draw_inverted == 0) ^ ((brush->flag & BRUSH_DIR_IN) == 0)) && + BKE_brush_sculpt_has_secondary_color(brush)) { outline_col = brush->sub_col; } diff --git a/source/blender/editors/sculpt_paint/paint_ops.c b/source/blender/editors/sculpt_paint/paint_ops.c index 4f6b3d100c5..2899cfeedcf 100644 --- a/source/blender/editors/sculpt_paint/paint_ops.c +++ b/source/blender/editors/sculpt_paint/paint_ops.c @@ -264,9 +264,14 @@ static int brush_reset_exec(bContext *C, wmOperator *UNUSED(op)) if (!ob || !brush) return OPERATOR_CANCELLED; - if (ob->mode & OB_MODE_SCULPT) - BKE_brush_sculpt_reset(brush); /* TODO: other modes */ + if (ob->mode & OB_MODE_SCULPT) { + BKE_brush_sculpt_reset(brush); + } + else { + return OPERATOR_CANCELLED; + } + WM_event_add_notifier(C, NC_BRUSH | NA_EDITED, brush); return OPERATOR_FINISHED; } diff --git a/source/blender/editors/sculpt_paint/paint_vertex.c b/source/blender/editors/sculpt_paint/paint_vertex.c index 5d87030bc67..fd88ea2d15f 100644 --- a/source/blender/editors/sculpt_paint/paint_vertex.c +++ b/source/blender/editors/sculpt_paint/paint_vertex.c @@ -105,6 +105,7 @@ struct NormalAnglePrecalc { static void view_angle_limits_init( struct NormalAnglePrecalc *a, float angle, bool do_mask_normal) { + angle = RAD2DEGF(angle); a->do_mask_normal = do_mask_normal; if (do_mask_normal) { a->angle_inner = angle; @@ -144,8 +145,8 @@ static float view_angle_limits_apply_falloff( static bool vwpaint_use_normal(const VPaint *vp) { - return ((vp->flag & VP_FLAG_PROJECT_BACKFACE) == 0) || - ((vp->flag & VP_FLAG_PROJECT_FLAT) == 0); + return ((vp->paint.brush->flag & BRUSH_FRONTFACE) != 0) || + ((vp->paint.brush->flag & BRUSH_FRONTFACE_FALLOFF) != 0); } @@ -357,7 +358,7 @@ static float calc_vp_alpha_col_dl( if (strength > 0.0f) { float alpha = brush_alpha_pressure * strength; - if ((vp->flag & VP_FLAG_PROJECT_FLAT) == 0) { + if ((vp->paint.brush->flag & BRUSH_FRONTFACE_FALLOFF) != 0) { float dvec[3]; /* transpose ! */ @@ -1390,7 +1391,8 @@ static bool wpaint_stroke_test_start(bContext *C, wmOperator *op, const float mo wpd = MEM_callocN(sizeof(struct WPaintData), "WPaintData"); paint_stroke_set_mode_data(stroke, wpd); view3d_set_viewcontext(C, &wpd->vc); - view_angle_limits_init(&wpd->normal_angle_precalc, vp->normal_angle, (vp->flag & VP_FLAG_PROJECT_FLAT) == 0); + view_angle_limits_init(&wpd->normal_angle_precalc, vp->paint.brush->falloff_angle, + (vp->paint.brush->flag & BRUSH_FRONTFACE_FALLOFF) != 0); wpd->active.index = vgroup_index.active; wpd->mirror.index = vgroup_index.mirror; @@ -1510,22 +1512,6 @@ static float wpaint_get_active_weight(const MDeformVert *dv, const WeightPaintIn } } -static SculptBrushTestFn sculpt_brush_test_init_with_falloff_shape( - SculptSession *ss, SculptBrushTest *test, char falloff_shape) -{ - sculpt_brush_test_init(ss, test); - SculptBrushTestFn sculpt_brush_test_sq_fn; - if (falloff_shape == VP_FALLOFF_SHAPE_SPHERE) { - sculpt_brush_test_sq_fn = sculpt_brush_test_sphere_sq; - } - else { - /* VP_FALLOFF_SHAPE_TUBE */ - plane_from_point_normal_v3(test->plane, test->location, ss->cache->view_normal); - sculpt_brush_test_sq_fn = sculpt_brush_test_circle_sq; - } - return sculpt_brush_test_sq_fn; -} - static void do_wpaint_brush_blur_task_cb_ex( void *userdata, void *UNUSED(userdata_chunk), const int n, const int UNUSED(thread_id)) { @@ -1546,7 +1532,9 @@ static void do_wpaint_brush_blur_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -1581,10 +1569,10 @@ static void do_wpaint_brush_blur_task_cb_ex( if (total_hit_loops != 0) { float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->wpd->normal_angle_precalc, angle_cos, &brush_strength))) { const float brush_fade = BKE_brush_curve_strength(brush, sqrtf(test.dist), cache->radius); @@ -1639,7 +1627,9 @@ static void do_wpaint_brush_smear_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -1657,10 +1647,10 @@ static void do_wpaint_brush_smear_task_cb_ex( if (!(use_face_sel || use_vert_sel) || mv_curr->flag & SELECT) { float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->wpd->normal_angle_precalc, angle_cos, &brush_strength))) { bool do_color = false; @@ -1739,7 +1729,9 @@ static void do_wpaint_brush_draw_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -1758,10 +1750,10 @@ static void do_wpaint_brush_draw_task_cb_ex( if (!(use_face_sel || use_vert_sel) || v_flag & SELECT) { float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->wpd->normal_angle_precalc, angle_cos, &brush_strength))) { const float brush_fade = BKE_brush_curve_strength(brush, sqrtf(test.dist), cache->radius); @@ -1804,7 +1796,9 @@ static void do_wpaint_brush_calc_average_weight_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -1813,7 +1807,7 @@ static void do_wpaint_brush_calc_average_weight_cb_ex( /* Test to see if the vertex coordinates are within the spherical brush region. */ if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; if (angle_cos > 0.0 && BKE_brush_curve_strength(data->brush, sqrtf(test.dist), cache->radius) > 0.0) { const int v_index = ccgdm ? data->me->mloop[vd.grid_indices[vd.g]].v : vd.vert_indices[vd.i]; // const float grid_alpha = ccgdm ? 1.0f / vd.gridsize : 1.0f; @@ -1906,7 +1900,7 @@ static PBVHNode **vwpaint_pbvh_gather_generic( PBVHNode **nodes = NULL; /* Build a list of all nodes that are potentially within the brush's area of influence */ - if (wp->falloff_shape == VP_FALLOFF_SHAPE_SPHERE) { + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_SPHERE) { SculptSearchSphereData data = { .ss = ss, .sd = sd, @@ -1923,7 +1917,7 @@ static PBVHNode **vwpaint_pbvh_gather_generic( } else { struct DistRayAABB_Precalc dist_ray_to_aabb_precalc; - dist_squared_ray_to_aabb_precalc(&dist_ray_to_aabb_precalc, ss->cache->location, ss->cache->view_normal); + dist_squared_ray_to_aabb_v3_precalc(&dist_ray_to_aabb_precalc, ss->cache->location, ss->cache->view_normal); SculptSearchCircleData data = { .ss = ss, .sd = sd, @@ -2404,7 +2398,8 @@ static bool vpaint_stroke_test_start(bContext *C, struct wmOperator *op, const f vpd = MEM_callocN(sizeof(*vpd), "VPaintData"); paint_stroke_set_mode_data(stroke, vpd); view3d_set_viewcontext(C, &vpd->vc); - view_angle_limits_init(&vpd->normal_angle_precalc, vp->normal_angle, (vp->flag & VP_FLAG_PROJECT_FLAT) == 0); + view_angle_limits_init(&vpd->normal_angle_precalc, vp->paint.brush->falloff_angle, + (vp->paint.brush->flag & BRUSH_FRONTFACE_FALLOFF) != 0); vpd->paintcol = vpaint_get_current_col(scene, vp); @@ -2476,7 +2471,7 @@ static void do_vpaint_brush_calc_average_color_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -2545,7 +2540,9 @@ static void do_vpaint_brush_draw_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -2566,10 +2563,10 @@ static void do_vpaint_brush_draw_task_cb_ex( * (ie splash prevention factor), and only paint front facing verts. */ float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->vpd->normal_angle_precalc, angle_cos, &brush_strength))) { const float brush_fade = BKE_brush_curve_strength(brush, sqrtf(test.dist), cache->radius); @@ -2634,7 +2631,9 @@ static void do_vpaint_brush_blur_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -2652,10 +2651,10 @@ static void do_vpaint_brush_blur_task_cb_ex( if (!use_vert_sel || mv->flag & SELECT) { float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->vpd->normal_angle_precalc, angle_cos, &brush_strength))) { const float brush_fade = BKE_brush_curve_strength(brush, sqrtf(test.dist), cache->radius); @@ -2747,7 +2746,9 @@ static void do_vpaint_brush_smear_task_cb_ex( SculptBrushTest test; SculptBrushTestFn sculpt_brush_test_sq_fn = - sculpt_brush_test_init_with_falloff_shape(ss, &test, data->vp->falloff_shape); + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + const float *sculpt_normal_frontface = + sculpt_brush_frontface_normal_from_falloff_shape(ss, data->brush->falloff_shape); /* For each vertex */ PBVHVertexIter vd; @@ -2767,10 +2768,10 @@ static void do_vpaint_brush_smear_task_cb_ex( * (ie splash prevention factor), and only paint front facing verts. */ float brush_strength = cache->bstrength; const float angle_cos = (use_normal && vd.no) ? - dot_vf3vs3(ss->cache->sculpt_normal_symm, vd.no) : 1.0f; - if (((data->vp->flag & VP_FLAG_PROJECT_BACKFACE) || + dot_vf3vs3(sculpt_normal_frontface, vd.no) : 1.0f; + if (((brush->flag & BRUSH_FRONTFACE) == 0 || (angle_cos > 0.0f)) && - ((data->vp->flag & VP_FLAG_PROJECT_FLAT) || + ((brush->flag & BRUSH_FRONTFACE_FALLOFF) == 0 || view_angle_limits_apply_falloff(&data->vpd->normal_angle_precalc, angle_cos, &brush_strength))) { const float brush_fade = BKE_brush_curve_strength(brush, sqrtf(test.dist), cache->radius); diff --git a/source/blender/editors/sculpt_paint/paint_vertex_color_utils.c b/source/blender/editors/sculpt_paint/paint_vertex_color_utils.c index 3b3e7297aec..a8045232e05 100644 --- a/source/blender/editors/sculpt_paint/paint_vertex_color_utils.c +++ b/source/blender/editors/sculpt_paint/paint_vertex_color_utils.c @@ -298,7 +298,7 @@ BLI_INLINE uint mcol_colordodge(uint col1, uint col2, int fac) cp[0] = (mfac * cp1[0] + temp * fac) / 255; temp = (cp2[1] == 255) ? 255 : min_ii((cp1[1] * 225) / (255 - cp2[1]), 255); cp[1] = (mfac * cp1[1] + temp * fac) / 255; - temp = (cp2[2] == 255) ? 255 : min_ii((cp1[2] * 225 )/ (255 - cp2[2]), 255); + temp = (cp2[2] == 255) ? 255 : min_ii((cp1[2] * 225) / (255 - cp2[2]), 255); cp[2] = (mfac * cp1[2] + temp * fac) / 255; temp = (cp2[3] == 255) ? 255 : min_ii((cp1[3] * 225) / (255 - cp2[3]), 255); cp[3] = (mfac * cp1[3] + temp * fac) / 255; diff --git a/source/blender/editors/sculpt_paint/sculpt.c b/source/blender/editors/sculpt_paint/sculpt.c index aa8e536a430..fb5192b8072 100644 --- a/source/blender/editors/sculpt_paint/sculpt.c +++ b/source/blender/editors/sculpt_paint/sculpt.c @@ -524,7 +524,8 @@ void sculpt_brush_test_init(SculptSession *ss, SculptBrushTest *test) test->dist = 0.0f; /* just for initialize */ /* Only for 2D projection. */ - zero_v4(test->plane); + zero_v4(test->plane_view); + zero_v4(test->plane_tool); test->mirror_symmetry_pass = ss->cache->mirror_symmetry_pass; @@ -590,7 +591,7 @@ bool sculpt_brush_test_sphere_fast(const SculptBrushTest *test, const float co[3 bool sculpt_brush_test_circle_sq(SculptBrushTest *test, const float co[3]) { float co_proj[3]; - closest_to_plane_normalized_v3(co_proj, test->plane, co); + closest_to_plane_normalized_v3(co_proj, test->plane_view, co); float distsq = len_squared_v3v3(co_proj, test->location); if (distsq <= test->radius_squared) { @@ -634,6 +635,35 @@ bool sculpt_brush_test_cube(SculptBrushTest *test, const float co[3], float loca } } +SculptBrushTestFn sculpt_brush_test_init_with_falloff_shape( + SculptSession *ss, SculptBrushTest *test, char falloff_shape) +{ + sculpt_brush_test_init(ss, test); + SculptBrushTestFn sculpt_brush_test_sq_fn; + if (falloff_shape == PAINT_FALLOFF_SHAPE_SPHERE) { + sculpt_brush_test_sq_fn = sculpt_brush_test_sphere_sq; + } + else { + /* PAINT_FALLOFF_SHAPE_TUBE */ + plane_from_point_normal_v3(test->plane_view, test->location, ss->cache->view_normal); + sculpt_brush_test_sq_fn = sculpt_brush_test_circle_sq; + } + return sculpt_brush_test_sq_fn; +} + +const float *sculpt_brush_frontface_normal_from_falloff_shape( + SculptSession *ss, char falloff_shape) +{ + if (falloff_shape == PAINT_FALLOFF_SHAPE_SPHERE) { + return ss->cache->sculpt_normal_symm; + } + else { + /* PAINT_FALLOFF_SHAPE_TUBE */ + return ss->cache->view_normal; + } +} + + static float frontface(const Brush *br, const float sculpt_normal[3], const short no[3], const float fno[3]) { @@ -772,7 +802,6 @@ static void calc_area_normal_and_center_task_cb(void *userdata, const int n) float (*area_cos)[3] = data->area_cos; PBVHVertexIter vd; - SculptBrushTest test; SculptUndoNode *unode = NULL; float private_co[2][3] = {{0.0f}}; @@ -784,7 +813,10 @@ static void calc_area_normal_and_center_task_cb(void *userdata, const int n) unode = sculpt_undo_push_node(data->ob, data->nodes[n], SCULPT_UNDO_COORDS); use_original = (unode->co || unode->bm_entry); } - sculpt_brush_test_init(ss, &test); + + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); /* when the mesh is edited we can't rely on original coords @@ -807,7 +839,7 @@ static void calc_area_normal_and_center_task_cb(void *userdata, const int n) closest_on_tri_to_point_v3(co, test.location, UNPACK3(co_tri)); - if (sculpt_brush_test_sphere_fast(&test, co)) { + if (sculpt_brush_test_sq_fn(&test, co)) { float no[3]; int flip_index; @@ -841,7 +873,7 @@ static void calc_area_normal_and_center_task_cb(void *userdata, const int n) co = vd.co; } - if (sculpt_brush_test_sphere_fast(&test, co)) { + if (sculpt_brush_test_sq_fn(&test, co)) { float no_buf[3]; const float *no; int flip_index; @@ -959,7 +991,7 @@ void sculpt_pbvh_calc_area_normal( /* Intentionally set 'sd' to NULL since this is used for vertex paint too. */ SculptThreadedTaskData data = { - .sd = NULL, .ob = ob, .nodes = nodes, .totnode = totnode, + .sd = NULL, .ob = ob, .brush = brush, .nodes = nodes, .totnode = totnode, .has_bm_orco = has_bm_orco, .area_cos = NULL, .area_nos = area_nos, .count = count, }; BLI_mutex_init(&data.mutex); @@ -999,7 +1031,7 @@ static void calc_area_normal_and_center( /* Intentionally set 'sd' to NULL since this is used for vertex paint too. */ SculptThreadedTaskData data = { - .sd = NULL, .ob = ob, .nodes = nodes, .totnode = totnode, + .sd = NULL, .ob = ob, .brush = brush, .nodes = nodes, .totnode = totnode, .has_bm_orco = has_bm_orco, .area_cos = area_cos, .area_nos = area_nos, .count = count, }; BLI_mutex_init(&data.mutex); @@ -1246,10 +1278,10 @@ bool sculpt_search_circle_cb(PBVHNode *node, void *data_v) BKE_pbvh_node_get_BB(node, bb_min, bb_min); float dummy_co[3], dummy_depth; - const float dist_sq = dist_squared_ray_to_aabb( + const float dist_sq = dist_squared_ray_to_aabb_v3( data->dist_ray_to_aabb_precalc, bb_min, bb_max, dummy_co, &dummy_depth); - return dist_sq < data->radius_squared; + return dist_sq < data->radius_squared || 1; } /* Handles clipping against a mirror modifier and SCULPT_LOCK axis flags */ @@ -1268,6 +1300,37 @@ static void sculpt_clip(Sculpt *sd, SculptSession *ss, float co[3], const float } } +static PBVHNode **sculpt_pbvh_gather_generic( + Object *ob, Sculpt *sd, const Brush *brush, bool use_original, float radius_scale, int *r_totnode) +{ + SculptSession *ss = ob->sculpt; + PBVHNode **nodes = NULL; + + /* Build a list of all nodes that are potentially within the brush's area of influence */ + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_SPHERE) { + SculptSearchSphereData data = { + .ss = ss, + .sd = sd, + .radius_squared = SQUARE(ss->cache->radius * radius_scale), + .original = use_original, + }; + BKE_pbvh_search_gather(ss->pbvh, sculpt_search_sphere_cb, &data, &nodes, r_totnode); + } + else { + struct DistRayAABB_Precalc dist_ray_to_aabb_precalc; + dist_squared_ray_to_aabb_v3_precalc(&dist_ray_to_aabb_precalc, ss->cache->location, ss->cache->view_normal); + SculptSearchCircleData data = { + .ss = ss, + .sd = sd, + .radius_squared = SQUARE(ss->cache->radius * radius_scale), + .original = use_original, + .dist_ray_to_aabb_precalc = &dist_ray_to_aabb_precalc, + }; + BKE_pbvh_search_gather(ss->pbvh, sculpt_search_circle_cb, &data, &nodes, r_totnode); + } + return nodes; +} + /* Calculate primary direction of movement for many brushes */ static void calc_sculpt_normal( Sculpt *sd, Object *ob, @@ -1314,6 +1377,10 @@ static void update_sculpt_normal(Sculpt *sd, Object *ob, (cache->first_time || !(brush->flag & BRUSH_ORIGINAL_NORMAL))) { calc_sculpt_normal(sd, ob, nodes, totnode, cache->sculpt_normal); + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(cache->sculpt_normal, cache->sculpt_normal, cache->view_normal); + normalize_v3(cache->sculpt_normal); + } copy_v3_v3(cache->sculpt_normal_symm, cache->sculpt_normal); } else { @@ -1538,18 +1605,26 @@ typedef struct { SculptSession *ss; const float *ray_start, *ray_normal; bool hit; - float dist; + float depth; bool original; - PBVHNode* node; } SculptRaycastData; typedef struct { const float *ray_start, *ray_normal; bool hit; - float dist; + float depth; float detail; } SculptDetailRaycastData; +typedef struct { + SculptSession *ss; + const float *ray_start, *ray_normal; + bool hit; + float depth; + float dist_sq_to_ray; + bool original; +} SculptFindNearestToRayData; + static void do_smooth_brush_mesh_task_cb_ex( void *userdata, void *UNUSED(userdata_chunk), const int n, const int thread_id) { @@ -1561,19 +1636,19 @@ static void do_smooth_brush_mesh_task_cb_ex( float bstrength = data->strength; PBVHVertexIter vd; - SculptBrushTest test; CLAMP(bstrength, 0.0f, 1.0f); - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, - smooth_mask ? 0.0f : (vd.mask ? *vd.mask : 0.0f), - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, smooth_mask ? 0.0f : (vd.mask ? *vd.mask : 0.0f), thread_id); if (smooth_mask) { float val = neighbor_average_mask(ss, vd.vert_indices[vd.i]) - *vd.mask; val *= fade * bstrength; @@ -1609,18 +1684,19 @@ static void do_smooth_brush_bmesh_task_cb_ex( float bstrength = data->strength; PBVHVertexIter vd; - SculptBrushTest test; CLAMP(bstrength, 0.0f, 1.0f); - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, smooth_mask ? 0.0f : *vd.mask, - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, smooth_mask ? 0.0f : *vd.mask, thread_id); if (smooth_mask) { float val = bmesh_neighbor_average_mask(vd.bm_vert, vd.cd_vert_mask_offset) - *vd.mask; val *= fade * bstrength; @@ -1656,7 +1732,6 @@ static void do_smooth_brush_multires_task_cb_ex( const bool smooth_mask = data->smooth_mask; float bstrength = data->strength; - SculptBrushTest test; CCGElem **griddata, *gddata; CCGKey key; @@ -1669,7 +1744,9 @@ static void do_smooth_brush_multires_task_cb_ex( int *grid_indices, totgrid, gridsize; int i, x, y; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); CLAMP(bstrength, 0.0f, 1.0f); @@ -1756,10 +1833,11 @@ static void do_smooth_brush_multires_task_cb_ex( fno = CCG_elem_offset_no(&key, gddata, index); mask = CCG_elem_offset_mask(&key, gddata, index); - if (sculpt_brush_test_sphere(&test, co)) { + if (sculpt_brush_test_sq_fn(&test, co)) { const float strength_mask = (smooth_mask ? 0.0f : *mask); const float fade = bstrength * tex_strength( - ss, brush, co, test.dist, NULL, fno, strength_mask, thread_id); + ss, brush, co, sqrtf(test.dist), + NULL, fno, strength_mask, thread_id); float f = 1.0f / 16.0f; if (x == 0 || x == gridsize - 1) @@ -1870,14 +1948,17 @@ static void do_mask_brush_draw_task_cb_ex( const float bstrength = ss->cache->bstrength; PBVHVertexIter vd; - SculptBrushTest test; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { - const float fade = tex_strength(ss, brush, vd.co, test.dist, vd.no, vd.fno, 0.0f, thread_id); + if (sculpt_brush_test_sq_fn(&test, vd.co)) { + const float fade = tex_strength( + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, 0.0f, thread_id); (*vd.mask) += fade * bstrength; CLAMP(*vd.mask, 0, 1); @@ -1927,19 +2008,21 @@ static void do_draw_brush_task_cb_ex( const float *offset = data->offset; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { /* offset vertex */ const float fade = tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], offset, fade); @@ -1977,6 +2060,9 @@ static void do_draw_brush(Sculpt *sd, Object *ob, PBVHNode **nodes, int totnode) ((sd->flags & SCULPT_USE_OPENMP) && totnode > SCULPT_THREADED_LIMIT), false); } +/** + * Used for 'SCULPT_TOOL_CREASE' and 'SCULPT_TOOL_BLOB' + */ static void do_crease_brush_task_cb_ex( void *userdata, void *UNUSED(userdata_chunk), const int n, const int thread_id) { @@ -1988,24 +2074,30 @@ static void do_crease_brush_task_cb_ex( const float *offset = data->offset; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { /* offset vertex */ const float fade = tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); float val1[3]; float val2[3]; /* first we pinch */ sub_v3_v3v3(val1, test.location, vd.co); + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(val1, val1, ss->cache->view_normal); + } + mul_v3_fl(val1, fade * flippedbstrength); sculpt_project_v3(spvc, val1, val1); @@ -2073,22 +2165,27 @@ static void do_pinch_brush_task_cb_ex( const Brush *brush = data->brush; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); float val[3]; sub_v3_v3v3(val, test.location, vd.co); + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(val, val, ss->cache->view_normal); + } mul_v3_v3fl(proxy[vd.i], val, fade); if (vd.mvert) @@ -2120,7 +2217,6 @@ static void do_grab_brush_task_cb_ex( const float *grab_delta = data->grab_delta; PBVHVertexIter vd; - SculptBrushTest test; SculptOrigVertData orig_data; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; @@ -2129,16 +2225,18 @@ static void do_grab_brush_task_cb_ex( proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { sculpt_orig_vert_data_update(&orig_data, &vd); - if (sculpt_brush_test_sphere(&test, orig_data.co)) { + if (sculpt_brush_test_sq_fn(&test, orig_data.co)) { const float fade = bstrength * tex_strength( - ss, brush, orig_data.co, test.dist, orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, orig_data.co, sqrtf(test.dist), + orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], grab_delta, fade); @@ -2180,19 +2278,21 @@ static void do_nudge_brush_task_cb_ex( const float *cono = data->cono; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], cono, fade); @@ -2235,7 +2335,6 @@ static void do_snake_hook_brush_task_cb_ex( const float *grab_delta = data->grab_delta; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; const bool do_rake_rotation = ss->cache->is_rake_rotation_valid; @@ -2245,13 +2344,16 @@ static void do_snake_hook_brush_task_cb_ex( proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], grab_delta, fade); @@ -2260,6 +2362,9 @@ static void do_snake_hook_brush_task_cb_ex( float delta_pinch_init[3], delta_pinch[3]; sub_v3_v3v3(delta_pinch, vd.co, test.location); + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(delta_pinch, delta_pinch, ss->cache->true_view_normal); + } /* important to calculate based on the grabbed location (intentionally ignore fade here). */ add_v3_v3(delta_pinch, grab_delta); @@ -2335,7 +2440,6 @@ static void do_thumb_brush_task_cb_ex( const float *cono = data->cono; PBVHVertexIter vd; - SculptBrushTest test; SculptOrigVertData orig_data; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; @@ -2344,16 +2448,18 @@ static void do_thumb_brush_task_cb_ex( proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { sculpt_orig_vert_data_update(&orig_data, &vd); - if (sculpt_brush_test_sphere(&test, orig_data.co)) { + if (sculpt_brush_test_sq_fn(&test, orig_data.co)) { const float fade = bstrength * tex_strength( - ss, brush, orig_data.co, test.dist, orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, orig_data.co, sqrtf(test.dist), + orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], cono, fade); @@ -2395,7 +2501,6 @@ static void do_rotate_brush_task_cb_ex( const float angle = data->angle; PBVHVertexIter vd; - SculptBrushTest test; SculptOrigVertData orig_data; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; @@ -2404,17 +2509,19 @@ static void do_rotate_brush_task_cb_ex( proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { sculpt_orig_vert_data_update(&orig_data, &vd); - if (sculpt_brush_test_sphere(&test, orig_data.co)) { + if (sculpt_brush_test_sq_fn(&test, orig_data.co)) { float vec[3], rot[3][3]; const float fade = bstrength * tex_strength( - ss, brush, orig_data.co, test.dist, orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, orig_data.co, sqrtf(test.dist), + orig_data.no, NULL, vd.mask ? *vd.mask : 0.0f, thread_id); sub_v3_v3v3(vec, orig_data.co, ss->cache->location); axis_angle_normalized_to_mat3(rot, ss->cache->sculpt_normal_symm, angle * fade); @@ -2457,7 +2564,6 @@ static void do_layer_brush_task_cb_ex( const float *offset = data->offset; PBVHVertexIter vd; - SculptBrushTest test; SculptOrigVertData orig_data; float *layer_disp; const float bstrength = ss->cache->bstrength; @@ -2472,15 +2578,18 @@ static void do_layer_brush_task_cb_ex( layer_disp = BKE_pbvh_node_layer_disp_get(ss->pbvh, data->nodes[n]); BLI_mutex_unlock(&data->mutex); - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { sculpt_orig_vert_data_update(&orig_data, &vd); - if (sculpt_brush_test_sphere(&test, orig_data.co)) { + if (sculpt_brush_test_sq_fn(&test, orig_data.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); float *disp = &layer_disp[vd.i]; float val[3]; @@ -2540,19 +2649,21 @@ static void do_inflate_brush_task_cb_ex( const Brush *brush = data->brush; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, test.dist, vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); float val[3]; if (vd.fno) @@ -2615,6 +2726,10 @@ static void calc_sculpt_plane( case SCULPT_DISP_DIR_AREA: calc_area_normal_and_center(sd, ob, nodes, totnode, r_area_no, r_area_co); + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(r_area_no, r_area_no, ss->cache->view_normal); + normalize_v3(r_area_no); + } break; default: @@ -2700,29 +2815,31 @@ static void do_flatten_brush_task_cb_ex( const float *area_co = data->area_co; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); - plane_from_point_normal_v3(test.plane, area_co, area_no); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + + plane_from_point_normal_v3(test.plane_tool, area_co, area_no); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere_sq(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { float intr[3]; float val[3]; - closest_to_plane_normalized_v3(intr, test.plane, vd.co); + closest_to_plane_normalized_v3(intr, test.plane_tool, vd.co); sub_v3_v3v3(val, intr, vd.co); if (plane_trim(ss->cache, brush, val)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, sqrtf(test.dist), vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], val, fade); @@ -2776,24 +2893,26 @@ static void do_clay_brush_task_cb_ex( const float *area_co = data->area_co; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const bool flip = (ss->cache->bstrength < 0); const float bstrength = flip ? -ss->cache->bstrength : ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); - plane_from_point_normal_v3(test.plane, area_co, area_no); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + + plane_from_point_normal_v3(test.plane_tool, area_co, area_no); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere_sq(&test, vd.co)) { - if (plane_point_side_flip(vd.co, test.plane, flip)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { + if (plane_point_side_flip(vd.co, test.plane_tool, flip)) { float intr[3]; float val[3]; - closest_to_plane_normalized_v3(intr, test.plane, vd.co); + closest_to_plane_normalized_v3(intr, test.plane_tool, vd.co); sub_v3_v3v3(val, intr, vd.co); @@ -2801,8 +2920,8 @@ static void do_clay_brush_task_cb_ex( /* note, the normal from the vertices is ignored, * causes glitch with planes, see: T44390 */ const float fade = bstrength * tex_strength( - ss, brush, vd.co, sqrtf(test.dist), vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], val, fade); @@ -2869,16 +2988,16 @@ static void do_clay_strips_brush_task_cb_ex( proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; sculpt_brush_test_init(ss, &test); - plane_from_point_normal_v3(test.plane, area_co, area_no_sp); + plane_from_point_normal_v3(test.plane_tool, area_co, area_no_sp); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { if (sculpt_brush_test_cube(&test, vd.co, mat)) { - if (plane_point_side_flip(vd.co, test.plane, flip)) { + if (plane_point_side_flip(vd.co, test.plane_tool, flip)) { float intr[3]; float val[3]; - closest_to_plane_normalized_v3(intr, test.plane, vd.co); + closest_to_plane_normalized_v3(intr, test.plane_tool, vd.co); sub_v3_v3v3(val, intr, vd.co); @@ -2886,8 +3005,8 @@ static void do_clay_strips_brush_task_cb_ex( /* note, the normal from the vertices is ignored, * causes glitch with planes, see: T44390 */ const float fade = bstrength * tex_strength( - ss, brush, vd.co, ss->cache->radius * test.dist, - vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, ss->cache->radius * test.dist, + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], val, fade); @@ -2970,30 +3089,32 @@ static void do_fill_brush_task_cb_ex( const float *area_co = data->area_co; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); - plane_from_point_normal_v3(test.plane, area_co, area_no); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + + plane_from_point_normal_v3(test.plane_tool, area_co, area_no); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere_sq(&test, vd.co)) { - if (plane_point_side(vd.co, test.plane)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { + if (plane_point_side(vd.co, test.plane_tool)) { float intr[3]; float val[3]; - closest_to_plane_normalized_v3(intr, test.plane, vd.co); + closest_to_plane_normalized_v3(intr, test.plane_tool, vd.co); sub_v3_v3v3(val, intr, vd.co); if (plane_trim(ss->cache, brush, val)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, sqrtf(test.dist), - vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], val, fade); @@ -3049,30 +3170,31 @@ static void do_scrape_brush_task_cb_ex( const float *area_co = data->area_co; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; const float bstrength = ss->cache->bstrength; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); - plane_from_point_normal_v3(test.plane, area_co, area_no); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); + plane_from_point_normal_v3(test.plane_tool, area_co, area_no); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere_sq(&test, vd.co)) { - if (!plane_point_side(vd.co, test.plane)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { + if (!plane_point_side(vd.co, test.plane_tool)) { float intr[3]; float val[3]; - closest_to_plane_normalized_v3(intr, test.plane, vd.co); + closest_to_plane_normalized_v3(intr, test.plane_tool, vd.co); sub_v3_v3v3(val, intr, vd.co); if (plane_trim(ss->cache, brush, val)) { const float fade = bstrength * tex_strength( - ss, brush, vd.co, sqrtf(test.dist), vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], val, fade); @@ -3127,18 +3249,19 @@ static void do_gravity_task_cb_ex( float *offset = data->offset; PBVHVertexIter vd; - SculptBrushTest test; float (*proxy)[3]; proxy = BKE_pbvh_node_add_proxy(ss->pbvh, data->nodes[n])->co; - sculpt_brush_test_init(ss, &test); + SculptBrushTest test; + SculptBrushTestFn sculpt_brush_test_sq_fn = + sculpt_brush_test_init_with_falloff_shape(ss, &test, data->brush->falloff_shape); BKE_pbvh_vertex_iter_begin(ss->pbvh, data->nodes[n], vd, PBVH_ITER_UNIQUE) { - if (sculpt_brush_test_sphere_sq(&test, vd.co)) { + if (sculpt_brush_test_sq_fn(&test, vd.co)) { const float fade = tex_strength( - ss, brush, vd.co, sqrtf(test.dist), vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, - thread_id); + ss, brush, vd.co, sqrtf(test.dist), + vd.no, vd.fno, vd.mask ? *vd.mask : 0.0f, thread_id); mul_v3_v3fl(proxy[vd.i], offset, fade); @@ -3222,22 +3345,12 @@ void sculpt_vertcos_to_key(Object *ob, KeyBlock *kb, float (*vertCos)[3]) static void sculpt_topology_update(Sculpt *sd, Object *ob, Brush *brush, UnifiedPaintSettings *UNUSED(ups)) { SculptSession *ss = ob->sculpt; - SculptSearchSphereData data; - PBVHNode **nodes = NULL; - float radius; - int n, totnode; - /* Build a list of all nodes that are potentially within the - * brush's area of influence */ - data.ss = ss; - data.sd = sd; - - radius = ss->cache->radius * 1.25f; - - data.radius_squared = radius * radius; - data.original = sculpt_tool_needs_original(brush->sculpt_tool) ? true : ss->cache->original; - - BKE_pbvh_search_gather(ss->pbvh, sculpt_search_sphere_cb, &data, &nodes, &totnode); + int n, totnode; + /* Build a list of all nodes that are potentially within the brush's area of influence */ + const bool use_original = sculpt_tool_needs_original(brush->sculpt_tool) ? true : ss->cache->original; + const float radius_scale = 1.25f; + PBVHNode **nodes = sculpt_pbvh_gather_generic(ob, sd, brush, use_original, radius_scale, &totnode); /* Only act if some verts are inside the brush area */ if (totnode) { @@ -3269,8 +3382,10 @@ static void sculpt_topology_update(Sculpt *sd, Object *ob, Brush *brush, Unified BKE_pbvh_bmesh_update_topology( ss->pbvh, mode, ss->cache->location, - (brush->flag & BRUSH_FRONTFACE) ? ss->cache->view_normal : NULL, - ss->cache->radius); + ss->cache->view_normal, + ss->cache->radius, + (brush->flag & BRUSH_FRONTFACE) != 0, + (brush->falloff_shape != PAINT_FALLOFF_SHAPE_SPHERE)); } MEM_freeN(nodes); @@ -3293,16 +3408,12 @@ static void do_brush_action_task_cb(void *userdata, const int n) static void do_brush_action(Sculpt *sd, Object *ob, Brush *brush, UnifiedPaintSettings *ups) { SculptSession *ss = ob->sculpt; - SculptSearchSphereData data; - PBVHNode **nodes = NULL; int totnode; /* Build a list of all nodes that are potentially within the brush's area of influence */ - data.ss = ss; - data.sd = sd; - data.radius_squared = ss->cache->radius_squared; - data.original = sculpt_tool_needs_original(brush->sculpt_tool) ? true : ss->cache->original; - BKE_pbvh_search_gather(ss->pbvh, sculpt_search_sphere_cb, &data, &nodes, &totnode); + const bool use_original = sculpt_tool_needs_original(brush->sculpt_tool) ? true : ss->cache->original; + const float radius_scale = 1.0f; + PBVHNode **nodes = sculpt_pbvh_gather_generic(ob, sd, brush, use_original, radius_scale, &totnode); /* Only act if some verts are inside the brush area */ if (totnode) { @@ -4120,7 +4231,6 @@ static void sculpt_update_brush_delta(UnifiedPaintSettings *ups, Object *ob, Bru sub_v3_v3v3(cache->grab_delta, grab_location, cache->old_grab_location); } - invert_m4_m4(imat, ob->obmat); mul_mat3_m4_v3(imat, cache->grab_delta); break; @@ -4130,6 +4240,10 @@ static void sculpt_update_brush_delta(UnifiedPaintSettings *ups, Object *ob, Bru zero_v3(cache->grab_delta); } + if (brush->falloff_shape == PAINT_FALLOFF_SHAPE_TUBE) { + project_plane_v3_v3v3(cache->grab_delta, cache->grab_delta, ss->cache->true_view_normal); + } + copy_v3_v3(cache->old_grab_location, grab_location); if (tool == SCULPT_TOOL_GRAB) @@ -4330,13 +4444,40 @@ static void sculpt_raycast_cb(PBVHNode *node, void *data_v, float *tmin) } if (BKE_pbvh_node_raycast(srd->ss->pbvh, node, origco, use_origco, - srd->ray_start, srd->ray_normal, &srd->dist)) + srd->ray_start, srd->ray_normal, &srd->depth)) { srd->hit = 1; - *tmin = srd->dist; + *tmin = srd->depth; + } + } +} + +static void sculpt_find_nearest_to_ray_cb(PBVHNode *node, void *data_v, float *tmin) +{ + if (BKE_pbvh_node_get_tmin(node) < *tmin) { + SculptFindNearestToRayData *srd = data_v; + float (*origco)[3] = NULL; + bool use_origco = false; + + if (srd->original && srd->ss->cache) { + if (BKE_pbvh_type(srd->ss->pbvh) == PBVH_BMESH) { + use_origco = true; + } + else { + /* intersect with coordinates from before we started stroke */ + SculptUndoNode *unode = sculpt_undo_get_node(node); + origco = (unode) ? unode->co : NULL; + use_origco = origco ? true : false; + } + } - //for vwpaint testing - srd->node = node; + if (BKE_pbvh_node_find_nearest_to_ray( + srd->ss->pbvh, node, origco, use_origco, + srd->ray_start, srd->ray_normal, + &srd->depth, &srd->dist_sq_to_ray)) + { + srd->hit = 1; + *tmin = srd->dist_sq_to_ray; } } } @@ -4346,10 +4487,10 @@ static void sculpt_raycast_detail_cb(PBVHNode *node, void *data_v, float *tmin) if (BKE_pbvh_node_get_tmin(node) < *tmin) { SculptDetailRaycastData *srd = data_v; if (BKE_pbvh_bmesh_node_raycast_detail(node, srd->ray_start, srd->ray_normal, - &srd->dist, &srd->detail)) + &srd->depth, &srd->detail)) { srd->hit = 1; - *tmin = srd->dist; + *tmin = srd->depth; } } } @@ -4396,8 +4537,7 @@ bool sculpt_stroke_get_location(bContext *C, float out[3], const float mouse[2]) Object *ob; SculptSession *ss; StrokeCache *cache; - float ray_start[3], ray_end[3], ray_normal[3], dist; - SculptRaycastData srd; + float ray_start[3], ray_end[3], ray_normal[3], depth; bool original; ViewContext vc; @@ -4411,28 +4551,58 @@ bool sculpt_stroke_get_location(bContext *C, float out[3], const float mouse[2]) sculpt_stroke_modifiers_check(C, ob); - dist = sculpt_raycast_init(&vc, mouse, ray_start, ray_end, ray_normal, original); + depth = sculpt_raycast_init(&vc, mouse, ray_start, ray_end, ray_normal, original); - srd.original = original; - srd.ss = ob->sculpt; - srd.hit = 0; - srd.ray_start = ray_start; - srd.ray_normal = ray_normal; - srd.dist = dist; - - BKE_pbvh_raycast(ss->pbvh, sculpt_raycast_cb, &srd, - ray_start, ray_normal, srd.original); + bool hit = false; + { + SculptRaycastData srd = { + .original = original, + .ss = ob->sculpt, + .hit = 0, + .ray_start = ray_start, + .ray_normal = ray_normal, + .depth = depth, + }; + BKE_pbvh_raycast( + ss->pbvh, sculpt_raycast_cb, &srd, + ray_start, ray_normal, srd.original); + if (srd.hit) { + hit = true; + copy_v3_v3(out, ray_normal); + mul_v3_fl(out, srd.depth); + add_v3_v3(out, ray_start); + } + } - copy_v3_v3(out, ray_normal); - mul_v3_fl(out, srd.dist); - add_v3_v3(out, ray_start); + if (hit == false) { + const Brush *brush = BKE_paint_brush(BKE_paint_get_active_from_context(C)); + if (ELEM(brush->falloff_shape, PAINT_FALLOFF_SHAPE_TUBE)) { + SculptFindNearestToRayData srd = { + .original = original, + .ss = ob->sculpt, + .hit = 0, + .ray_start = ray_start, + .ray_normal = ray_normal, + .depth = FLT_MAX, + .dist_sq_to_ray = FLT_MAX, + }; + BKE_pbvh_find_nearest_to_ray( + ss->pbvh, sculpt_find_nearest_to_ray_cb, &srd, + ray_start, ray_normal, srd.original); + if (srd.hit) { + hit = true; + copy_v3_v3(out, ray_normal); + mul_v3_fl(out, srd.depth); + add_v3_v3(out, ray_start); + } + } + } - //used in vwpaint - if (cache && srd.hit){ + if (cache && hit) { copy_v3_v3(cache->true_location, out); } - return srd.hit; + return hit; } static void sculpt_brush_init_tex(const Scene *scene, Sculpt *sd, SculptSession *ss) @@ -5475,7 +5645,7 @@ static int sculpt_detail_flood_fill_exec(bContext *C, wmOperator *UNUSED(op)) while (BKE_pbvh_bmesh_update_topology( ss->pbvh, PBVH_Collapse | PBVH_Subdivide, - center, NULL, size)) + center, NULL, size, false, false)) { for (i = 0; i < totnodes; i++) BKE_pbvh_node_mark_topology_update(nodes[i]); @@ -5511,7 +5681,7 @@ static void sample_detail(bContext *C, int ss_co[2]) ViewContext vc; Object *ob; Sculpt *sd; - float ray_start[3], ray_end[3], ray_normal[3], dist; + float ray_start[3], ray_end[3], ray_normal[3], depth; SculptDetailRaycastData srd; float mouse[2] = {ss_co[0], ss_co[1]}; view3d_set_viewcontext(C, &vc); @@ -5521,12 +5691,12 @@ static void sample_detail(bContext *C, int ss_co[2]) sculpt_stroke_modifiers_check(C, ob); - dist = sculpt_raycast_init(&vc, mouse, ray_start, ray_end, ray_normal, false); + depth = sculpt_raycast_init(&vc, mouse, ray_start, ray_end, ray_normal, false); srd.hit = 0; srd.ray_start = ray_start; srd.ray_normal = ray_normal; - srd.dist = dist; + srd.depth = depth; srd.detail = sd->constant_detail; BKE_pbvh_raycast(ob->sculpt->pbvh, sculpt_raycast_detail_cb, &srd, diff --git a/source/blender/editors/sculpt_paint/sculpt_intern.h b/source/blender/editors/sculpt_paint/sculpt_intern.h index b9b662f4157..aaea13ce5d0 100644 --- a/source/blender/editors/sculpt_paint/sculpt_intern.h +++ b/source/blender/editors/sculpt_paint/sculpt_intern.h @@ -182,7 +182,10 @@ typedef struct SculptBrushTest { int mirror_symmetry_pass; /* For circle (not sphere) projection. */ - float plane[4]; + float plane_view[4]; + + /* Some tool code uses a plane for it's calculateions. */ + float plane_tool[4]; /* View3d clipping - only set rv3d for clipping */ struct RegionView3D *clip_rv3d; @@ -213,6 +216,12 @@ bool sculpt_brush_test_cube(SculptBrushTest *test, const float co[3], float loca bool sculpt_brush_test_circle_sq(SculptBrushTest *test, const float co[3]); bool sculpt_search_sphere_cb(PBVHNode *node, void *data_v); bool sculpt_search_circle_cb(PBVHNode *node, void *data_v); + +SculptBrushTestFn sculpt_brush_test_init_with_falloff_shape( + SculptSession *ss, SculptBrushTest *test, char falloff_shape); +const float *sculpt_brush_frontface_normal_from_falloff_shape( + SculptSession *ss, char falloff_shape); + float tex_strength( struct SculptSession *ss, const struct Brush *br, const float point[3], diff --git a/source/blender/editors/space_file/file_ops.c b/source/blender/editors/space_file/file_ops.c index b7228f634bf..ab1bcbaa8b5 100644 --- a/source/blender/editors/space_file/file_ops.c +++ b/source/blender/editors/space_file/file_ops.c @@ -1302,7 +1302,6 @@ void file_sfile_filepath_set(SpaceFile *sfile, const char *filepath) if (BLI_is_dir(filepath)) { BLI_strncpy(sfile->params->dir, filepath, sizeof(sfile->params->dir)); - sfile->params->file[0] = '\0'; } else { if ((sfile->params->flag & FILE_DIRSEL_ONLY) == 0) { diff --git a/source/blender/editors/space_file/filelist.c b/source/blender/editors/space_file/filelist.c index 48f3dd4fde6..d670af59bc9 100644 --- a/source/blender/editors/space_file/filelist.c +++ b/source/blender/editors/space_file/filelist.c @@ -1994,8 +1994,7 @@ int ED_path_extension_type(const char *path) else if (BLI_testextensie(path, ".abc")) { return FILE_TYPE_ALEMBIC; } - else if (BLI_testextensie_array(path, imb_ext_image)) - { + else if (BLI_testextensie_array(path, imb_ext_image)) { return FILE_TYPE_IMAGE; } else if (BLI_testextensie(path, ".ogg")) { diff --git a/source/blender/editors/space_image/image_ops.c b/source/blender/editors/space_image/image_ops.c index 1f4848636be..1dae5aeafa5 100644 --- a/source/blender/editors/space_image/image_ops.c +++ b/source/blender/editors/space_image/image_ops.c @@ -1333,7 +1333,8 @@ static int image_open_exec(bContext *C, wmOperator *op) iuser->framenr = 1; if (ima->source == IMA_SRC_MOVIE) { iuser->offset = 0; - } else { + } + else { iuser->offset = frame_ofs - 1; } iuser->fie_ima = 2; diff --git a/source/blender/editors/space_node/node_templates.c b/source/blender/editors/space_node/node_templates.c index ec525e684b0..5d0877a1eff 100644 --- a/source/blender/editors/space_node/node_templates.c +++ b/source/blender/editors/space_node/node_templates.c @@ -683,10 +683,11 @@ static void ui_node_draw_input(uiLayout *layout, bContext *C, bNodeTree *ntree, RNA_pointer_create(&ntree->id, &RNA_Node, node, &nodeptr); /* indented label */ - for (i = 0; i < indent; i++) + for (i = 0; i < indent; i++) { label[i] = ' '; + } label[indent] = '\0'; - BLI_snprintf(label, UI_MAX_NAME_STR, "%s%s:", label, IFACE_(input->name)); + BLI_snprintf(label + indent, UI_MAX_NAME_STR - indent, "%s:", IFACE_(input->name)); /* split in label and value */ split = uiLayoutSplit(layout, 0.35f, false); diff --git a/source/blender/editors/space_view3d/view3d_ruler.c b/source/blender/editors/space_view3d/view3d_ruler.c index 7b7d996c77e..60e90121ad1 100644 --- a/source/blender/editors/space_view3d/view3d_ruler.c +++ b/source/blender/editors/space_view3d/view3d_ruler.c @@ -1041,9 +1041,12 @@ static int view3d_ruler_modal(bContext *C, wmOperator *op, const wmEvent *event) } case RETKEY: { - view3d_ruler_to_gpencil(C, ruler_info); - do_draw = true; - exit_code = OPERATOR_FINISHED; + /* Enter may be used to invoke from search. */ + if (event->val == KM_PRESS) { + view3d_ruler_to_gpencil(C, ruler_info); + do_draw = true; + exit_code = OPERATOR_FINISHED; + } break; } case DELKEY: diff --git a/source/blender/makesdna/DNA_brush_types.h b/source/blender/makesdna/DNA_brush_types.h index 67a35d4e207..c285b44c939 100644 --- a/source/blender/makesdna/DNA_brush_types.h +++ b/source/blender/makesdna/DNA_brush_types.h @@ -95,8 +95,12 @@ typedef struct Brush { float plane_offset; /* offset for plane brushes (clay, flatten, fill, scrape) */ int gradient_spacing; - int gradient_stroke_mode; /* source for stroke color gradient application */ - int gradient_fill_mode; /* source for fill tool color gradient application */ + char gradient_stroke_mode; /* source for stroke color gradient application */ + char gradient_fill_mode; /* source for fill tool color gradient application */ + + char pad; + char falloff_shape; /* Projection shape (sphere, circle) */ + float falloff_angle; char sculpt_tool; /* active sculpt tool */ char vertexpaint_tool; /* active vertex/weight paint blend mode (poorly named) */ @@ -181,13 +185,13 @@ typedef enum BrushGradientSourceFill { /* Brush.flag */ typedef enum BrushFlags { BRUSH_AIRBRUSH = (1 << 0), -// BRUSH_TORUS = (1 << 1), deprecated, use paint->symmetry_flags & PAINT_TILE_* + BRUSH_FLAG_DEPRECATED_1 = (1 << 1), BRUSH_ALPHA_PRESSURE = (1 << 2), BRUSH_SIZE_PRESSURE = (1 << 3), BRUSH_JITTER_PRESSURE = (1 << 4), BRUSH_SPACING_PRESSURE = (1 << 5), - BRUSH_UNUSED = (1 << 6), -// BRUSH_RAKE = (1 << 7), deprecated, use brush_angle_mode + BRUSH_FLAG_DEPRECATED_2 = (1 << 6), + BRUSH_FLAG_DEPRECATED_3 = (1 << 7), BRUSH_ANCHORED = (1 << 8), BRUSH_DIR_IN = (1 << 9), BRUSH_SPACE = (1 << 10), @@ -197,6 +201,7 @@ typedef enum BrushFlags { BRUSH_LOCK_ALPHA = (1 << 14), BRUSH_ORIGINAL_NORMAL = (1 << 15), BRUSH_OFFSET_PRESSURE = (1 << 16), + BRUSH_FLAG_DEPRECATED_4 = (1 << 17), BRUSH_SPACE_ATTEN = (1 << 18), BRUSH_ADAPTIVE_SPACE = (1 << 19), BRUSH_LOCK_SIZE = (1 << 20), @@ -204,7 +209,7 @@ typedef enum BrushFlags { BRUSH_EDGE_TO_EDGE = (1 << 22), BRUSH_DRAG_DOT = (1 << 23), BRUSH_INVERSE_SMOOTH_PRESSURE = (1 << 24), -// BRUSH_RANDOM_ROTATION = (1 << 25), deprecated, use brush_angle_mode + BRUSH_FRONTFACE_FALLOFF = (1 << 25), BRUSH_PLANE_TRIM = (1 << 26), BRUSH_FRONTFACE = (1 << 27), BRUSH_CUSTOM_ICON = (1 << 28), @@ -343,6 +348,12 @@ typedef enum BlurKernelType { KERNEL_BOX } BlurKernelType; +/* Brush.falloff_shape */ +enum { + PAINT_FALLOFF_SHAPE_SPHERE = 0, + PAINT_FALLOFF_SHAPE_TUBE = 1, +}; + #define MAX_BRUSH_PIXEL_RADIUS 500 #endif /* __DNA_BRUSH_TYPES_H__ */ diff --git a/source/blender/makesdna/DNA_scene_types.h b/source/blender/makesdna/DNA_scene_types.h index 85f2f6a9816..d81a4625144 100644 --- a/source/blender/makesdna/DNA_scene_types.h +++ b/source/blender/makesdna/DNA_scene_types.h @@ -1121,27 +1121,17 @@ typedef struct UvSculpt { /* Vertex Paint */ typedef struct VPaint { Paint paint; - short flag; - char falloff_shape, normal_angle; + char flag; + char pad[3]; int radial_symm[3]; /* For mirrored painting */ } VPaint; /* VPaint.flag */ enum { - VP_FLAG_PROJECT_BACKFACE = (1 << 0), - /* TODO */ - // VP_FLAG_PROJECT_XRAY = (1 << 1), - VP_FLAG_PROJECT_FLAT = (1 << 3), /* weight paint only */ VP_FLAG_VGROUP_RESTRICT = (1 << 7) }; -/* VPaint.falloff_shape */ -enum { - VP_FALLOFF_SHAPE_SPHERE = 0, - VP_FALLOFF_SHAPE_TUBE = 1, -}; - /* ------------------------------------------- */ /* GPencil Stroke Sculpting */ diff --git a/source/blender/makesrna/intern/rna_brush.c b/source/blender/makesrna/intern/rna_brush.c index 352706fd9b8..f79ff4486dd 100644 --- a/source/blender/makesrna/intern/rna_brush.c +++ b/source/blender/makesrna/intern/rna_brush.c @@ -238,11 +238,7 @@ static int rna_SculptToolCapabilities_has_sculpt_plane_get(PointerRNA *ptr) static int rna_SculptToolCapabilities_has_secondary_color_get(PointerRNA *ptr) { Brush *br = (Brush *)ptr->data; - return ELEM(br->sculpt_tool, - SCULPT_TOOL_BLOB, SCULPT_TOOL_CLAY, SCULPT_TOOL_CLAY_STRIPS, - SCULPT_TOOL_CREASE, SCULPT_TOOL_DRAW, SCULPT_TOOL_FILL, - SCULPT_TOOL_FLATTEN, SCULPT_TOOL_INFLATE, SCULPT_TOOL_PINCH, - SCULPT_TOOL_SCRAPE); + return BKE_brush_sculpt_has_secondary_color(br); } static int rna_SculptToolCapabilities_has_smooth_stroke_get(PointerRNA *ptr) @@ -1136,7 +1132,20 @@ static void rna_def_brush(BlenderRNA *brna) RNA_def_property_ui_text(prop, "Blur Mode", ""); RNA_def_property_update(prop, 0, "rna_Brush_update"); + prop = RNA_def_property(srna, "falloff_angle", PROP_FLOAT, PROP_ANGLE); + RNA_def_property_float_sdna(prop, NULL, "falloff_angle"); + RNA_def_property_range(prop, 0, M_PI / 2); + RNA_def_property_ui_text(prop, "Falloff Angle", + "Paint most on faces pointing towards the view according to this angle"); + RNA_def_property_update(prop, 0, "rna_Brush_update"); + /* flag */ + /* This is an enum but its unlikely we add other shapes, so expose as a boolean. */ + prop = RNA_def_property(srna, "use_projected", PROP_BOOLEAN, PROP_NONE); + RNA_def_property_boolean_sdna(prop, NULL, "falloff_shape", BRUSH_AIRBRUSH); + RNA_def_property_ui_text(prop, "2D Falloff", "Apply brush influence in 2D circle instead of a sphere"); + RNA_def_property_update(prop, 0, "rna_Brush_update"); + prop = RNA_def_property(srna, "use_airbrush", PROP_BOOLEAN, PROP_NONE); RNA_def_property_boolean_sdna(prop, NULL, "flag", BRUSH_AIRBRUSH); RNA_def_property_ui_text(prop, "Airbrush", "Keep applying paint effect while holding mouse (spray)"); @@ -1213,6 +1222,11 @@ static void rna_def_brush(BlenderRNA *brna) RNA_def_property_ui_text(prop, "Use Front-Face", "Brush only affects vertexes that face the viewer"); RNA_def_property_update(prop, 0, "rna_Brush_update"); + prop = RNA_def_property(srna, "use_frontface_falloff", PROP_BOOLEAN, PROP_NONE); + RNA_def_property_boolean_sdna(prop, NULL, "flag", BRUSH_FRONTFACE_FALLOFF); + RNA_def_property_ui_text(prop, "Use Front-Face Falloff", "Blend brush influence by how much they face the front"); + RNA_def_property_update(prop, 0, "rna_Brush_update"); + prop = RNA_def_property(srna, "use_anchor", PROP_BOOLEAN, PROP_NONE); RNA_def_property_boolean_sdna(prop, NULL, "flag", BRUSH_ANCHORED); RNA_def_property_ui_text(prop, "Anchored", "Keep the brush anchored to the initial location"); diff --git a/source/blender/makesrna/intern/rna_sculpt_paint.c b/source/blender/makesrna/intern/rna_sculpt_paint.c index f2c4e19a6e4..cfe8670b0f5 100644 --- a/source/blender/makesrna/intern/rna_sculpt_paint.c +++ b/source/blender/makesrna/intern/rna_sculpt_paint.c @@ -690,38 +690,12 @@ static void rna_def_vertex_paint(BlenderRNA *brna) RNA_def_struct_path_func(srna, "rna_VertexPaint_path"); RNA_def_struct_ui_text(srna, "Vertex Paint", "Properties of vertex and weight paint mode"); - static EnumPropertyItem prop_falloff_items[] = { - {VP_FALLOFF_SHAPE_SPHERE, "SPHERE", 0, "Sphere", "Spherical falloff from the brush"}, - {VP_FALLOFF_SHAPE_TUBE, "TUBE", 0, "Circle", "Circular falloff from the brush along the view"}, - {0, NULL, 0, NULL, NULL} - }; - - prop = RNA_def_property(srna, "falloff_shape", PROP_ENUM, PROP_NONE); - RNA_def_property_enum_sdna(prop, NULL, "falloff_shape"); - RNA_def_property_enum_items(prop, prop_falloff_items); - RNA_def_property_ui_text(prop, "Falloff", ""); - RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL); - - prop = RNA_def_property(srna, "use_backface_culling", PROP_BOOLEAN, PROP_NONE); - RNA_def_property_boolean_negative_sdna(prop, NULL, "flag", VP_FLAG_PROJECT_BACKFACE); - RNA_def_property_ui_text(prop, "Cull", "Ignore vertices pointing away from the view (faster)"); - RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL); - - prop = RNA_def_property(srna, "use_normal_falloff", PROP_BOOLEAN, PROP_NONE); - RNA_def_property_boolean_negative_sdna(prop, NULL, "flag", VP_FLAG_PROJECT_FLAT); - RNA_def_property_ui_text(prop, "Normals", "Paint most on faces pointing towards the view"); - RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL); - /* weight paint only */ prop = RNA_def_property(srna, "use_group_restrict", PROP_BOOLEAN, PROP_NONE); RNA_def_property_boolean_sdna(prop, NULL, "flag", VP_FLAG_VGROUP_RESTRICT); RNA_def_property_ui_text(prop, "Restrict", "Restrict painting to vertices in the group"); RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL); - prop = RNA_def_property(srna, "normal_angle", PROP_INT, PROP_UNSIGNED); - RNA_def_property_range(prop, 0, 90); - RNA_def_property_ui_text(prop, "Angle", "Paint most on faces pointing towards the view according to this angle"); - /* Mirroring */ prop = RNA_def_property(srna, "radial_symmetry", PROP_INT, PROP_XYZ); RNA_def_property_int_sdna(prop, NULL, "radial_symm"); diff --git a/source/blender/python/generic/bgl.c b/source/blender/python/generic/bgl.c index c290e5bd7a0..5c938ef2ee3 100644 --- a/source/blender/python/generic/bgl.c +++ b/source/blender/python/generic/bgl.c @@ -806,7 +806,7 @@ static PyObject *Buffer_new(PyTypeObject *UNUSED(type), PyObject *args, PyObject "`GL_TYPE` and `typestr` of object with buffer interface do not match. '%s'", pybuffer.format); } else if (ndimensions != pybuffer.ndim || - !compare_dimensions(ndimensions, dimensions, pybuffer.shape)) + !compare_dimensions(ndimensions, dimensions, pybuffer.shape)) { PyErr_Format(PyExc_TypeError, "array size does not match"); } diff --git a/source/blender/python/intern/bpy.c b/source/blender/python/intern/bpy.c index 4a29d4f8da1..c576786a54c 100644 --- a/source/blender/python/intern/bpy.c +++ b/source/blender/python/intern/bpy.c @@ -119,10 +119,11 @@ static PyObject *bpy_blend_paths(PyObject *UNUSED(self), PyObject *args, PyObjec bool absolute = false; bool packed = false; bool local = false; - static const char *kwlist[] = {"absolute", "packed", "local", NULL}; - if (!PyArg_ParseTupleAndKeywords( - args, kw, "|O&O&O&:blend_paths", (char **)kwlist, + static const char *_keywords[] = {"absolute", "packed", "local", NULL}; + static _PyArg_Parser _parser = {"|O&O&O&:blend_paths", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, PyC_ParseBool, &absolute, PyC_ParseBool, &packed, PyC_ParseBool, &local)) @@ -148,13 +149,18 @@ static PyObject *bpy_user_resource(PyObject *UNUSED(self), PyObject *args, PyObj const char *type; const char *subdir = NULL; int folder_id; - static const char *kwlist[] = {"type", "subdir", NULL}; const char *path; - if (!PyArg_ParseTupleAndKeywords(args, kw, "s|s:user_resource", (char **)kwlist, &type, &subdir)) + static const char *_keywords[] = {"type", "subdir", NULL}; + static _PyArg_Parser _parser = {"s|s:user_resource", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, + &type, &subdir)) + { return NULL; - + } + /* stupid string compare */ if (STREQ(type, "DATAFILES")) folder_id = BLENDER_USER_DATAFILES; else if (STREQ(type, "CONFIG")) folder_id = BLENDER_USER_CONFIG; @@ -192,12 +198,17 @@ static PyObject *bpy_resource_path(PyObject *UNUSED(self), PyObject *args, PyObj { const char *type; int major = BLENDER_VERSION / 100, minor = BLENDER_VERSION % 100; - static const char *kwlist[] = {"type", "major", "minor", NULL}; int folder_id; const char *path; - if (!PyArg_ParseTupleAndKeywords(args, kw, "s|ii:resource_path", (char **)kwlist, &type, &major, &minor)) + static const char *_keywords[] = {"type", "major", "minor", NULL}; + static _PyArg_Parser _parser = {"s|ii:resource_path", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, + &type, &major, &minor)) + { return NULL; + } /* stupid string compare */ if (STREQ(type, "USER")) folder_id = BLENDER_RESOURCE_PATH_USER; diff --git a/source/blender/python/intern/bpy_library_load.c b/source/blender/python/intern/bpy_library_load.c index 5e8cb6968b0..79081e47e27 100644 --- a/source/blender/python/intern/bpy_library_load.c +++ b/source/blender/python/intern/bpy_library_load.c @@ -183,17 +183,17 @@ PyDoc_STRVAR(bpy_lib_load_doc, " :arg relative: When True the path is stored relative to the open blend file.\n" " :type relative: bool\n" ); -static PyObject *bpy_lib_load(PyObject *UNUSED(self), PyObject *args, PyObject *kwds) +static PyObject *bpy_lib_load(PyObject *UNUSED(self), PyObject *args, PyObject *kw) { - static const char *kwlist[] = {"filepath", "link", "relative", NULL}; Main *bmain = CTX_data_main(BPy_GetContext()); BPy_Library *ret; const char *filename = NULL; bool is_rel = false, is_link = false; - if (!PyArg_ParseTupleAndKeywords( - args, kwds, - "s|O&O&:load", (char **)kwlist, + static const char *_keywords[] = {"filepath", "link", "relative", NULL}; + static _PyArg_Parser _parser = {"s|O&O&:load", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, &filename, PyC_ParseBool, &is_link, PyC_ParseBool, &is_rel)) diff --git a/source/blender/python/intern/bpy_library_write.c b/source/blender/python/intern/bpy_library_write.c index bf91253141a..c054183034a 100644 --- a/source/blender/python/intern/bpy_library_write.c +++ b/source/blender/python/intern/bpy_library_write.c @@ -69,24 +69,23 @@ PyDoc_STRVAR(bpy_lib_write_doc, " :arg compress: When True, write a compressed blend file.\n" " :type compress: bool\n" ); -static PyObject *bpy_lib_write(PyObject *UNUSED(self), PyObject *args, PyObject *kwds) +static PyObject *bpy_lib_write(PyObject *UNUSED(self), PyObject *args, PyObject *kw) { - static const char *kwlist[] = { - "filepath", "datablocks", - /* optional */ - "relative_remap", "fake_user", "compress", - NULL, - }; - /* args */ const char *filepath; char filepath_abs[FILE_MAX]; PyObject *datablocks = NULL; bool use_relative_remap = false, use_fake_user = false, use_compress = false; - if (!PyArg_ParseTupleAndKeywords( - args, kwds, - "sO!|$O&O&O&:write", (char **)kwlist, + static const char *_keywords[] = { + "filepath", "datablocks", + /* optional */ + "relative_remap", "fake_user", "compress", + NULL, + }; + static _PyArg_Parser _parser = {"sO!|$O&O&O&:write", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, &filepath, &PySet_Type, &datablocks, PyC_ParseBool, &use_relative_remap, diff --git a/source/blender/python/intern/bpy_operator.c b/source/blender/python/intern/bpy_operator.c index bd3e5736c6c..c1fcb0792af 100644 --- a/source/blender/python/intern/bpy_operator.c +++ b/source/blender/python/intern/bpy_operator.c @@ -450,7 +450,7 @@ static PyObject *pyop_getinstance(PyObject *UNUSED(self), PyObject *value) op = PyMem_MALLOC(sizeof(wmOperator)); memset(op, 0, sizeof(wmOperator)); #endif - BLI_strncpy(op->idname, op->idname, sizeof(op->idname)); /* in case its needed */ + BLI_strncpy(op->idname, ot->idname, sizeof(op->idname)); /* in case its needed */ op->type = ot; RNA_pointer_create(NULL, &RNA_Operator, op, &ptr); diff --git a/source/blender/python/intern/bpy_rna_id_collection.c b/source/blender/python/intern/bpy_rna_id_collection.c index 72705ffb3fb..8def52dc8fb 100644 --- a/source/blender/python/intern/bpy_rna_id_collection.c +++ b/source/blender/python/intern/bpy_rna_id_collection.c @@ -163,7 +163,6 @@ static PyObject *bpy_user_map(PyObject *UNUSED(self), PyObject *args, PyObject * Main *bmain = G.main; /* XXX Ugly, but should work! */ #endif - static const char *kwlist[] = {"subset", "key_types", "value_types", NULL}; PyObject *subset = NULL; PyObject *key_types = NULL; @@ -173,9 +172,10 @@ static PyObject *bpy_user_map(PyObject *UNUSED(self), PyObject *args, PyObject * PyObject *ret = NULL; - - if (!PyArg_ParseTupleAndKeywords( - args, kwds, "|O$O!O!:user_map", (char **)kwlist, + static const char *_keywords[] = {"subset", "key_types", "value_types", NULL}; + static _PyArg_Parser _parser = {"|O$O!O!:user_map", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kwds, &_parser, &subset, &PySet_Type, &key_types, &PySet_Type, &val_types)) diff --git a/source/blender/python/intern/bpy_utils_units.c b/source/blender/python/intern/bpy_utils_units.c index e0c70483e3c..0ef689d1a5a 100644 --- a/source/blender/python/intern/bpy_utils_units.c +++ b/source/blender/python/intern/bpy_utils_units.c @@ -174,8 +174,6 @@ PyDoc_STRVAR(bpyunits_to_value_doc, ); static PyObject *bpyunits_to_value(PyObject *UNUSED(self), PyObject *args, PyObject *kw) { - static const char *kwlist[] = {"unit_system", "unit_category", "str_input", "str_ref_unit", NULL}; - char *usys_str = NULL, *ucat_str = NULL, *inpt = NULL, *uref = NULL; const float scale = 1.0f; @@ -185,8 +183,13 @@ static PyObject *bpyunits_to_value(PyObject *UNUSED(self), PyObject *args, PyObj int usys, ucat; PyObject *ret; - if (!PyArg_ParseTupleAndKeywords(args, kw, "sss#|z:bpy.utils.units.to_value", (char **)kwlist, - &usys_str, &ucat_str, &inpt, &str_len, &uref)) + static const char *_keywords[] = { + "unit_system", "unit_category", "str_input", "str_ref_unit", NULL, + }; + static _PyArg_Parser _parser = {"sss#|z:to_value", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, + &usys_str, &ucat_str, &inpt, &str_len, &uref)) { return NULL; } @@ -244,9 +247,6 @@ PyDoc_STRVAR(bpyunits_to_string_doc, ); static PyObject *bpyunits_to_string(PyObject *UNUSED(self), PyObject *args, PyObject *kw) { - static const char *kwlist[] = {"unit_system", "unit_category", "value", - "precision", "split_unit", "compatible_unit", NULL}; - char *usys_str = NULL, *ucat_str = NULL; double value = 0.0; int precision = 3; @@ -254,9 +254,13 @@ static PyObject *bpyunits_to_string(PyObject *UNUSED(self), PyObject *args, PyOb int usys, ucat; - if (!PyArg_ParseTupleAndKeywords( - args, kw, - "ssd|iO&O&:bpy.utils.units.to_string", (char **)kwlist, + static const char *_keywords[] = { + "unit_system", "unit_category", "value", + "precision", "split_unit", "compatible_unit", NULL, + }; + static _PyArg_Parser _parser = {"ssd|iO&O&:to_string", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, &usys_str, &ucat_str, &value, &precision, PyC_ParseBool, &split_unit, PyC_ParseBool, &compatible_unit)) diff --git a/source/blender/python/intern/gpu.c b/source/blender/python/intern/gpu.c index d3453c94450..a7a0ae78f26 100644 --- a/source/blender/python/intern/gpu.c +++ b/source/blender/python/intern/gpu.c @@ -207,7 +207,7 @@ PyDoc_STRVAR(GPU_export_shader_doc, " :return: Dictionary defining the shader, uniforms and attributes.\n" " :rtype: Dict" ); -static PyObject *GPU_export_shader(PyObject *UNUSED(self), PyObject *args, PyObject *kwds) +static PyObject *GPU_export_shader(PyObject *UNUSED(self), PyObject *args, PyObject *kw) { PyObject *pyscene; PyObject *pymat; @@ -224,11 +224,14 @@ static PyObject *GPU_export_shader(PyObject *UNUSED(self), PyObject *args, PyObj GPUInputUniform *uniform; GPUInputAttribute *attribute; - static const char *kwlist[] = {"scene", "material", NULL}; - - if (!PyArg_ParseTupleAndKeywords(args, kwds, "OO:export_shader", (char **)(kwlist), &pyscene, &pymat)) + static const char *_keywords[] = {"scene", "material", NULL}; + static _PyArg_Parser _parser = {"OO:export_shader", _keywords, 0}; + if (!_PyArg_ParseTupleAndKeywordsFast( + args, kw, &_parser, + &pyscene, &pymat)) + { return NULL; - + } scene = (Scene *)PyC_RNA_AsPointer(pyscene, "Scene"); if (scene == NULL) { return NULL; diff --git a/source/blender/render/intern/source/convertblender.c b/source/blender/render/intern/source/convertblender.c index 9d2ac76f7e6..1329ad484ea 100644 --- a/source/blender/render/intern/source/convertblender.c +++ b/source/blender/render/intern/source/convertblender.c @@ -272,7 +272,7 @@ static void calc_tangent_vector(ObjectRen *obr, VlakRen *vlr, int do_tangent) } else return; - tangent_from_uv(uv1, uv2, uv3, v1->co, v2->co, v3->co, vlr->n, tang); + tangent_from_uv_v3(uv1, uv2, uv3, v1->co, v2->co, v3->co, vlr->n, tang); if (do_tangent) { tav= RE_vertren_get_tangent(obr, v1, 1); @@ -284,7 +284,7 @@ static void calc_tangent_vector(ObjectRen *obr, VlakRen *vlr, int do_tangent) } if (v4) { - tangent_from_uv(uv1, uv3, uv4, v1->co, v3->co, v4->co, vlr->n, tang); + tangent_from_uv_v3(uv1, uv3, uv4, v1->co, v3->co, v4->co, vlr->n, tang); if (do_tangent) { tav= RE_vertren_get_tangent(obr, v1, 1); @@ -399,7 +399,7 @@ static void calc_vertexnormals(Render *UNUSED(re), ObjectRen *obr, bool do_verte float *n4= (vlr->v4)? vlr->v4->n: NULL; const float *c4= (vlr->v4)? vlr->v4->co: NULL; - accumulate_vertex_normals(vlr->v1->n, vlr->v2->n, vlr->v3->n, n4, + accumulate_vertex_normals_v3(vlr->v1->n, vlr->v2->n, vlr->v3->n, n4, vlr->n, vlr->v1->co, vlr->v2->co, vlr->v3->co, c4); } if (do_tangent) { diff --git a/source/blender/windowmanager/intern/wm_window.c b/source/blender/windowmanager/intern/wm_window.c index efa0932ba92..a3a01b36ec1 100644 --- a/source/blender/windowmanager/intern/wm_window.c +++ b/source/blender/windowmanager/intern/wm_window.c @@ -516,8 +516,12 @@ static void wm_window_ghostwindow_add(wmWindowManager *wm, const char *title, wm /* store actual window size in blender window */ bounds = GHOST_GetClientBounds(win->ghostwin); - win->sizex = GHOST_GetWidthRectangle(bounds); - win->sizey = GHOST_GetHeightRectangle(bounds); + + /* win32: gives undefined window size when minimized */ + if (GHOST_GetWindowState(win->ghostwin) != GHOST_kWindowStateMinimized) { + win->sizex = GHOST_GetWidthRectangle(bounds); + win->sizey = GHOST_GetHeightRectangle(bounds); + } GHOST_DisposeRectangle(bounds); #ifndef __APPLE__ |