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

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

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