[Bf-blender-cvs] [122743e7f2] cycles_split_kernel: Cycles: Faster version of driver workaround for OpenCL
Mai Lavelle
noreply at git.blender.org
Sat Feb 18 11:34:27 CET 2017
Commit: 122743e7f23f8df908bcd10322d39f951ebc83e2
Author: Mai Lavelle
Date: Fri Feb 17 22:42:58 2017 -0500
Branches: cycles_split_kernel
https://developer.blender.org/rB122743e7f23f8df908bcd10322d39f951ebc83e2
Cycles: Faster version of driver workaround for OpenCL
===================================================================
M intern/cycles/device/opencl/opencl.h
===================================================================
diff --git a/intern/cycles/device/opencl/opencl.h b/intern/cycles/device/opencl/opencl.h
index 74256ca9d8..6dab30d49c 100644
--- a/intern/cycles/device/opencl/opencl.h
+++ b/intern/cycles/device/opencl/opencl.h
@@ -29,24 +29,28 @@ CCL_NAMESPACE_BEGIN
/* Define CYCLES_DISABLE_DRIVER_WORKAROUNDS to disable workaounds for testing */
#ifndef CYCLES_DISABLE_DRIVER_WORKAROUNDS
-/* Work around AMD driver hangs by ensuring each command is finished before doing anything else. */
+/* Work around AMD driver hangs by setting an event callback which seems to make commands
+ * finish completely for some reason. This is faster than doing clFinish after every call.
+ */
+static cl_event driver_workaround_event;
+static void driver_workaround_event_callback_handler(cl_event event, cl_int, void*) {
+ clReleaseEvent(event);
+}
+
# undef clEnqueueNDRangeKernel
# define clEnqueueNDRangeKernel(a, b, c, d, e, f, g, h, i) \
- clFinish(a); \
- CLEW_GET_FUN(__clewEnqueueNDRangeKernel)(a, b, c, d, e, f, g, h, i); \
- clFinish(a);
+ CLEW_GET_FUN(__clewEnqueueNDRangeKernel)(a, b, c, d, e, f, g, h, &driver_workaround_event); \
+ clSetEventCallback(driver_workaround_event, CL_SUBMITTED, driver_workaround_event_callback_handler, NULL);
# undef clEnqueueWriteBuffer
# define clEnqueueWriteBuffer(a, b, c, d, e, f, g, h, i) \
- clFinish(a); \
- CLEW_GET_FUN(__clewEnqueueWriteBuffer)(a, b, c, d, e, f, g, h, i); \
- clFinish(a);
+ CLEW_GET_FUN(__clewEnqueueWriteBuffer)(a, b, c, d, e, f, g, h, &driver_workaround_event); \
+ clSetEventCallback(driver_workaround_event, CL_SUBMITTED, driver_workaround_event_callback_handler, NULL);
# undef clEnqueueReadBuffer
# define clEnqueueReadBuffer(a, b, c, d, e, f, g, h, i) \
- clFinish(a); \
- CLEW_GET_FUN(__clewEnqueueReadBuffer)(a, b, c, d, e, f, g, h, i); \
- clFinish(a);
+ CLEW_GET_FUN(__clewEnqueueReadBuffer)(a, b, c, d, e, f, g, h, &driver_workaround_event); \
+ clSetEventCallback(driver_workaround_event, CL_SUBMITTED, driver_workaround_event_callback_handler, NULL);
#endif /* CYCLES_DISABLE_DRIVER_WORKAROUNDS */
#define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
More information about the Bf-blender-cvs
mailing list