diff options
author | Philipp Oeser <info@graphics-engineer.com> | 2020-08-28 15:00:51 +0300 |
---|---|---|
committer | Philipp Oeser <info@graphics-engineer.com> | 2020-08-28 15:00:51 +0300 |
commit | 2bb60db94a25e69dd8527c3523a8ac1afdd0520d (patch) | |
tree | c1c49775967ee8b8524bd118e91848d75fd41665 /intern | |
parent | 0e021414fe9921164bd96340be8d9e2c49e4669a (diff) | |
parent | d3f2037966882d677711568c81de8b9b7eaec486 (diff) |
Merge branch 'blender-v2.90-release'
Diffstat (limited to 'intern')
-rw-r--r-- | intern/cycles/device/opencl/device_opencl_impl.cpp | 8 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_bake.cl | 14 |
2 files changed, 12 insertions, 10 deletions
diff --git a/intern/cycles/device/opencl/device_opencl_impl.cpp b/intern/cycles/device/opencl/device_opencl_impl.cpp index e851749949d..f0683d12f1f 100644 --- a/intern/cycles/device/opencl/device_opencl_impl.cpp +++ b/intern/cycles/device/opencl/device_opencl_impl.cpp @@ -864,6 +864,11 @@ void OpenCLDevice::load_preview_kernels() bool OpenCLDevice::wait_for_availability(const DeviceRequestedFeatures &requested_features) { + if (requested_features.use_baking) { + /* For baking, kernels have already been loaded in load_required_kernels(). */ + return true; + } + if (background) { load_kernel_task_pool.wait_work(); use_preview_kernels = false; @@ -1933,13 +1938,12 @@ void OpenCLDevice::bake(DeviceTask &task, RenderTile &rtile) kernel_set_args(kernel, start_arg_index, sample); enqueue_kernel(kernel, d_w, d_h); + clFinish(cqCommandQueue); rtile.sample = sample + 1; task.update_progress(&rtile, rtile.w * rtile.h); } - - clFinish(cqCommandQueue); } static bool kernel_build_opencl_2(cl_device_id cdDevice) diff --git a/intern/cycles/kernel/kernels/opencl/kernel_bake.cl b/intern/cycles/kernel/kernels/opencl/kernel_bake.cl index 041312b53cb..7b81e387467 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_bake.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_bake.cl @@ -12,12 +12,11 @@ __kernel void kernel_ocl_bake( ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, + ccl_global float *buffer, KERNEL_BUFFER_PARAMS, - int type, int filter, int sx, int sw, int offset, int sample) + int sx, int sy, int sw, int sh, int offset, int stride, int sample) { KernelGlobals kglobals, *kg = &kglobals; @@ -27,12 +26,11 @@ __kernel void kernel_ocl_bake( kernel_set_buffer_info(kg); int x = sx + ccl_global_id(0); + int y = sy + ccl_global_id(1); - if(x < sx + sw) { -#ifdef __NO_BAKING__ - output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); -#else - kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, x, offset, sample); + if(x < sx + sw && y < sy + sh) { +#ifndef __NO_BAKING__ + kernel_bake_evaluate(kg, buffer, sample, x, y, offset, stride); #endif } } |