diff options
author | Mai Lavelle <mai.lavelle@gmail.com> | 2017-02-07 09:03:17 +0300 |
---|---|---|
committer | Mai Lavelle <mai.lavelle@gmail.com> | 2017-02-09 05:56:23 +0300 |
commit | 5609ba7d78e2a16960638f5045d4dca6c89b6fd0 (patch) | |
tree | f6ed84d5b90a49803bc3e99b877ac7ea78e2684e /intern/cycles/device | |
parent | bea6cfec1a08a967db923b3534daf2b9329184c7 (diff) |
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.
Diffstat (limited to 'intern/cycles/device')
-rw-r--r-- | intern/cycles/device/opencl/opencl_base.cpp | 65 |
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); + } + } } } |