ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Fix T80149: Cycles OpenCL baking broken after changes to uses tiles for baking

Browse source 
We forgot to update this code as part of D3108. I'd like to include this in 2.90,
it's entirely broken now so can't really get any worse.

Differential Revision: https://developer.blender.org/D8738
This commit is contained in:
Brecht Van Lommel 2020-08-27 19:20:02 +02:00
parent 1725e46cee
commit d3f2037966
2 changed files with 12 additions and 10 deletions
Download patch file Download diff file Expand all files Collapse all files

8
intern/cycles/device/opencl/device_opencl_impl.cpp
View file

@ -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)

14
intern/cycles/kernel/kernels/opencl/kernel_bake.cl
View file

@ -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
}
}