Cycles: Fix OpenCL mem_zero and preview renders

Viewport rendering wasn't working, this was caused by two problems in mem_zero:

- `mem_zero` was being called before kernels were loaded so the `zero_buffer`
  kernel was unavailable. Added a fallback for this case.
- The global size used for enqueuing the `zero_buffer` kernel was large and
  causing the kernel to fail on some devices.

1 files changed, 44 insertions, 21 deletions

diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp
index 924641c124d..5a87808922c 100644
--- a/intern/cycles/device/opencl/opencl_base.cpp
+++ b/intern/cycles/device/opencl/opencl_base.cpp
@@ -320,38 +320,61 @@ void OpenCLDeviceBase::mem_copy_from(device_memory& mem, int y, int w, int h, in
 void OpenCLDeviceBase::mem_zero(device_memory& mem)
 {
 	if(mem.device_pointer) {
-		cl_kernel ckZeroBuffer = base_program(ustring("zero_buffer"));
+		if(base_program.is_loaded()) {
+			cl_kernel ckZeroBuffer = base_program(ustring("zero_buffer"));
 
-		size_t max_work_items[3];
-		clGetDeviceInfo(cdDevice, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, max_work_items, NULL);
-		size_t global_size = max_work_items[0] * max_work_items[1] * max_work_items[2];
+			size_t global_size[] = {256, 256};
+			size_t num_threads = global_size[0] * global_size[1];
 
-		cl_mem d_buffer = CL_MEM_PTR(mem.device_pointer);
-		unsigned long long d_offset = 0;
-		unsigned long long d_size = 0;
+			cl_mem d_buffer = CL_MEM_PTR(mem.device_pointer);
+			unsigned long long d_offset = 0;
+			unsigned long long d_size = 0;
 
-		while(d_offset + d_size < mem.memory_size()) {
-			d_size = std::min<unsigned long long>(global_size*sizeof(float4), mem.memory_size() - d_offset);
+			while(d_offset + d_size < mem.memory_size()) {
+				d_size = std::min<unsigned long long>(num_threads*sizeof(float4), mem.memory_size() - d_offset);
 
-			kernel_set_args(ckZeroBuffer, 0, d_buffer, d_size, d_offset);
+				kernel_set_args(ckZeroBuffer, 0, d_buffer, d_size, d_offset);
 
-			ciErr = clEnqueueNDRangeKernel(cqCommandQueue,
-				                           ckZeroBuffer,
-				                           1,
-				                           NULL,
-				                           &global_size,
-				                           NULL,
-				                           0,
-				                           NULL,
-				                           NULL);
-			opencl_assert_err(ciErr, "clEnqueueNDRangeKernel");
+				ciErr = clEnqueueNDRangeKernel(cqCommandQueue,
+						                       ckZeroBuffer,
+						                       2,
+						                       NULL,
+						                       global_size,
+						                       NULL,
+						                       0,
+						                       NULL,
+						                       NULL);
+				opencl_assert_err(ciErr, "clEnqueueNDRangeKernel");
 
-			d_offset += d_size;
+				d_offset += d_size;
+			}
 		}
 
 		if(mem.data_pointer) {
 			memset((void*)mem.data_pointer, 0, mem.memory_size());
 		}
+
+		if(!base_program.is_loaded()) {
+			void* zero = (void*)mem.data_pointer;
+
+			if(!mem.data_pointer) {
+				zero = util_aligned_malloc(mem.memory_size(), 16);
+				memset(zero, 0, mem.memory_size());
+			}
+
+			opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
+                                   CL_MEM_PTR(mem.device_pointer),
+                                   CL_TRUE,
+                                   0,
+                                   mem.memory_size(),
+                                   zero,
+                                   0,
+                                   NULL, NULL));
+
+			if(!mem.data_pointer) {
+				util_aligned_free(zero);
+			}
+		}
 	}
 }