Merge branch 'blender-v2.90-release'

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