Cycles: Split device_opencl.cpp into multiple files for easier maintenance
authorLukas Stockner <lukas.stockner@freenet.de>
Wed, 14 Sep 2016 21:47:54 +0000 (23:47 +0200)
committerLukas Stockner <lukas.stockner@freenet.de>
Sun, 9 Oct 2016 13:49:50 +0000 (15:49 +0200)
There are no user-visible changes, just some internal restructuring.

Differential Revision: https://developer.blender.org/D2231

intern/cycles/device/CMakeLists.txt
intern/cycles/device/device_opencl.cpp
intern/cycles/device/opencl/opencl.h [new file with mode: 0644]
intern/cycles/device/opencl/opencl_base.cpp [new file with mode: 0644]
intern/cycles/device/opencl/opencl_mega.cpp [new file with mode: 0644]
intern/cycles/device/opencl/opencl_split.cpp [new file with mode: 0644]
intern/cycles/device/opencl/opencl_util.cpp [new file with mode: 0644]

index c34677e1b78a0876b43e484457043c0bb958c429..966ff5e52ba409295e8069bb9153d67e744a05b1 100644 (file)
@@ -36,6 +36,15 @@ set(SRC
        device_task.cpp
 )
 
+set(SRC_OPENCL
+       opencl/opencl.h
+
+       opencl/opencl_base.cpp
+       opencl/opencl_mega.cpp
+       opencl/opencl_split.cpp
+       opencl/opencl_util.cpp
+)
+
 if(WITH_CYCLES_NETWORK)
        list(APPEND SRC
                device_network.cpp
@@ -67,4 +76,4 @@ endif()
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_device ${SRC} ${SRC_HEADERS})
+add_library(cycles_device ${SRC} ${SRC_OPENCL} ${SRC_HEADERS})
index 830e4d056b58a17b33add587c2e0c8da5b6c64bd..45cf6b074e9db119640936f88f6ea35aaddaf4b5 100644 (file)
 
 #ifdef WITH_OPENCL
 
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
+#include "opencl/opencl.h"
 
-#include "clew.h"
-
-#include "device.h"
 #include "device_intern.h"
 
-#include "buffers.h"
-
-#include "util_debug.h"
 #include "util_foreach.h"
 #include "util_logging.h"
-#include "util_map.h"
-#include "util_math.h"
-#include "util_md5.h"
-#include "util_opengl.h"
-#include "util_path.h"
-#include "util_time.h"
 
 CCL_NAMESPACE_BEGIN
 
-#define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
-
-/* Macro declarations used with split kernel */
-
-/* Macro to enable/disable work-stealing */
-#define __WORK_STEALING__
-
-#define SPLIT_KERNEL_LOCAL_SIZE_X 64
-#define SPLIT_KERNEL_LOCAL_SIZE_Y 1
-
-/* This value may be tuned according to the scene we are rendering.
- *
- * Modifying PATH_ITER_INC_FACTOR value proportional to number of expected
- * ray-bounces will improve performance.
- */
-#define PATH_ITER_INC_FACTOR 8
-
-/* When allocate global memory in chunks. We may not be able to
- * allocate exactly "CL_DEVICE_MAX_MEM_ALLOC_SIZE" bytes in chunks;
- * Since some bytes may be needed for aligning chunks of memory;
- * This is the amount of memory that we dedicate for that purpose.
- */
-#define DATA_ALLOCATION_MEM_FACTOR 5000000 //5MB
-
-struct OpenCLPlatformDevice {
-       OpenCLPlatformDevice(cl_platform_id platform_id,
-                            const string& platform_name,
-                            cl_device_id device_id,
-                            cl_device_type device_type,
-                            const string& device_name)
-         : platform_id(platform_id),
-           platform_name(platform_name),
-           device_id(device_id),
-           device_type(device_type),
-           device_name(device_name) {}
-       cl_platform_id platform_id;
-       string platform_name;
-       cl_device_id device_id;
-       cl_device_type device_type;
-       string device_name;
-};
-
-namespace {
-
-cl_device_type opencl_device_type()
-{
-       switch(DebugFlags().opencl.device_type)
-       {
-               case DebugFlags::OpenCL::DEVICE_NONE:
-                       return 0;
-               case DebugFlags::OpenCL::DEVICE_ALL:
-                       return CL_DEVICE_TYPE_ALL;
-               case DebugFlags::OpenCL::DEVICE_DEFAULT:
-                       return CL_DEVICE_TYPE_DEFAULT;
-               case DebugFlags::OpenCL::DEVICE_CPU:
-                       return CL_DEVICE_TYPE_CPU;
-               case DebugFlags::OpenCL::DEVICE_GPU:
-                       return CL_DEVICE_TYPE_GPU;
-               case DebugFlags::OpenCL::DEVICE_ACCELERATOR:
-                       return CL_DEVICE_TYPE_ACCELERATOR;
-               default:
-                       return CL_DEVICE_TYPE_ALL;
-       }
-}
-
-inline bool opencl_kernel_use_debug()
-{
-       return DebugFlags().opencl.debug;
-}
-
-bool opencl_kernel_use_advanced_shading(const string& platform)
-{
-       /* keep this in sync with kernel_types.h! */
-       if(platform == "NVIDIA CUDA")
-               return true;
-       else if(platform == "Apple")
-               return true;
-       else if(platform == "AMD Accelerated Parallel Processing")
-               return true;
-       else if(platform == "Intel(R) OpenCL")
-               return true;
-       /* Make sure officially unsupported OpenCL platforms
-        * does not set up to use advanced shading.
-        */
-       return false;
-}
-
-bool opencl_kernel_use_split(const string& platform_name,
-                             const cl_device_type device_type)
-{
-       if(DebugFlags().opencl.kernel_type == DebugFlags::OpenCL::KERNEL_SPLIT) {
-               VLOG(1) << "Forcing split kernel to use.";
-               return true;
-       }
-       if(DebugFlags().opencl.kernel_type == DebugFlags::OpenCL::KERNEL_MEGA) {
-               VLOG(1) << "Forcing mega kernel to use.";
-               return false;
-       }
-       /* TODO(sergey): Replace string lookups with more enum-like API,
-        * similar to device/vendor checks blender's gpu.
-        */
-       if(platform_name == "AMD Accelerated Parallel Processing" &&
-          device_type == CL_DEVICE_TYPE_GPU)
-       {
-               return true;
-       }
-       return false;
-}
-
-bool opencl_device_supported(const string& platform_name,
-                             const cl_device_id device_id)
-{
-       cl_device_type device_type;
-       clGetDeviceInfo(device_id,
-                       CL_DEVICE_TYPE,
-                       sizeof(cl_device_type),
-                       &device_type,
-                       NULL);
-       if(platform_name == "AMD Accelerated Parallel Processing" &&
-          device_type == CL_DEVICE_TYPE_GPU)
-       {
-               return true;
-       }
-       if(platform_name == "Apple" && device_type == CL_DEVICE_TYPE_GPU) {
-               return true;
-       }
-       return false;
-}
-
-bool opencl_platform_version_check(cl_platform_id platform,
-                                   string *error = NULL)
-{
-       const int req_major = 1, req_minor = 1;
-       int major, minor;
-       char version[256];
-       clGetPlatformInfo(platform,
-                         CL_PLATFORM_VERSION,
-                         sizeof(version),
-                         &version,
-                         NULL);
-       if(sscanf(version, "OpenCL %d.%d", &major, &minor) < 2) {
-               if(error != NULL) {
-                       *error = string_printf("OpenCL: failed to parse platform version string (%s).", version);
-               }
-               return false;
-       }
-       if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
-               if(error != NULL) {
-                       *error = string_printf("OpenCL: platform version 1.1 or later required, found %d.%d", major, minor);
-               }
-               return false;
-       }
-       if(error != NULL) {
-               *error = "";
-       }
-       return true;
-}
-
-bool opencl_device_version_check(cl_device_id device,
-                                 string *error = NULL)
-{
-       const int req_major = 1, req_minor = 1;
-       int major, minor;
-       char version[256];
-       clGetDeviceInfo(device,
-                       CL_DEVICE_OPENCL_C_VERSION,
-                       sizeof(version),
-                       &version,
-                       NULL);
-       if(sscanf(version, "OpenCL C %d.%d", &major, &minor) < 2) {
-               if(error != NULL) {
-                       *error = string_printf("OpenCL: failed to parse OpenCL C version string (%s).", version);
-               }
-               return false;
-       }
-       if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
-               if(error != NULL) {
-                       *error = string_printf("OpenCL: C version 1.1 or later required, found %d.%d", major, minor);
-               }
-               return false;
-       }
-       if(error != NULL) {
-               *error = "";
-       }
-       return true;
-}
-
-void opencl_get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices)
-{
-       const bool force_all_platforms =
-               (DebugFlags().opencl.kernel_type != DebugFlags::OpenCL::KERNEL_DEFAULT);
-       const cl_device_type device_type = opencl_device_type();
-       static bool first_time = true;
-#define FIRST_VLOG(severity) if(first_time) VLOG(severity)
-
-       usable_devices->clear();
-
-       if(device_type == 0) {
-               FIRST_VLOG(2) << "OpenCL devices are forced to be disabled.";
-               first_time = false;
-               return;
-       }
-
-       vector<cl_device_id> device_ids;
-       cl_uint num_devices = 0;
-       vector<cl_platform_id> platform_ids;
-       cl_uint num_platforms = 0;
-
-       /* Get devices. */
-       if(clGetPlatformIDs(0, NULL, &num_platforms) != CL_SUCCESS ||
-          num_platforms == 0)
-       {
-               FIRST_VLOG(2) << "No OpenCL platforms were found.";
-               first_time = false;
-               return;
-       }
-       platform_ids.resize(num_platforms);
-       if(clGetPlatformIDs(num_platforms, &platform_ids[0], NULL) != CL_SUCCESS) {
-               FIRST_VLOG(2) << "Failed to fetch platform IDs from the driver..";
-               first_time = false;
-               return;
-       }
-       /* Devices are numbered consecutively across platforms. */
-       for(int platform = 0; platform < num_platforms; platform++) {
-               cl_platform_id platform_id = platform_ids[platform];
-               char pname[256];
-               if(clGetPlatformInfo(platform_id,
-                                    CL_PLATFORM_NAME,
-                                    sizeof(pname),
-                                    &pname,
-                                    NULL) != CL_SUCCESS)
-               {
-                       FIRST_VLOG(2) << "Failed to get platform name, ignoring.";
-                       continue;
-               }
-               string platform_name = pname;
-               FIRST_VLOG(2) << "Enumerating devices for platform "
-                             << platform_name << ".";
-               if(!opencl_platform_version_check(platform_id)) {
-                       FIRST_VLOG(2) << "Ignoring platform " << platform_name
-                                     << " due to too old compiler version.";
-                       continue;
-               }
-               num_devices = 0;
-               if(clGetDeviceIDs(platform_id,
-                                 device_type,
-                                 0,
-                                 NULL,
-                                 &num_devices) != CL_SUCCESS || num_devices == 0)
-               {
-                       FIRST_VLOG(2) << "Ignoring platform " << platform_name
-                                     << ", failed to fetch number of devices.";
-                       continue;
-               }
-               device_ids.resize(num_devices);
-               if(clGetDeviceIDs(platform_id,
-                                 device_type,
-                                 num_devices,
-                                 &device_ids[0],
-                                 NULL) != CL_SUCCESS)
-               {
-                       FIRST_VLOG(2) << "Ignoring platform " << platform_name
-                                     << ", failed to fetch devices list.";
-                       continue;
-               }
-               for(int num = 0; num < num_devices; num++) {
-                       cl_device_id device_id = device_ids[num];
-                       char device_name[1024] = "\0";
-                       if(clGetDeviceInfo(device_id,
-                                          CL_DEVICE_NAME,
-                                          sizeof(device_name),
-                                          &device_name,
-                                          NULL) != CL_SUCCESS)
-                       {
-                               FIRST_VLOG(2) << "Failed to fetch device name, ignoring.";
-                               continue;
-                       }
-                       if(!opencl_device_version_check(device_id)) {
-                               FIRST_VLOG(2) << "Ignoring device " << device_name
-                                             << " due to old compiler version.";
-                               continue;
-                       }
-                       if(force_all_platforms ||
-                          opencl_device_supported(platform_name, device_id))
-                       {
-                               cl_device_type device_type;
-                               if(clGetDeviceInfo(device_id,
-                                                  CL_DEVICE_TYPE,
-                                                  sizeof(cl_device_type),
-                                                  &device_type,
-                                                  NULL) != CL_SUCCESS)
-                               {
-                                       FIRST_VLOG(2) << "Ignoring device " << device_name
-                                                     << ", failed to fetch device type.";
-                                       continue;
-                               }
-                               FIRST_VLOG(2) << "Adding new device " << device_name << ".";
-                               usable_devices->push_back(OpenCLPlatformDevice(platform_id,
-                                                                              platform_name,
-                                                                              device_id,
-                                                                              device_type,
-                                                                              device_name));
-                       }
-                       else {
-                               FIRST_VLOG(2) << "Ignoring device " << device_name
-                                             << ", not officially supported yet.";
-                       }
-               }
-       }
-       first_time = false;
-}
-
-}  /* namespace */
-
-/* Thread safe cache for contexts and programs.
- *
- * TODO(sergey): Make it more generous, so it can contain any type of program
- * without hardcoding possible program types in the slot.
- */
-class OpenCLCache
-{
-       struct Slot
-       {
-               thread_mutex *mutex;
-               cl_context context;
-               /* cl_program for shader, bake, film_convert kernels (used in OpenCLDeviceBase) */
-               cl_program ocl_dev_base_program;
-               /* cl_program for megakernel (used in OpenCLDeviceMegaKernel) */
-               cl_program ocl_dev_megakernel_program;
-
-               Slot() : mutex(NULL),
-                        context(NULL),
-                        ocl_dev_base_program(NULL),
-                        ocl_dev_megakernel_program(NULL) {}
-
-               Slot(const Slot& rhs)
-                   : mutex(rhs.mutex),
-                     context(rhs.context),
-                     ocl_dev_base_program(rhs.ocl_dev_base_program),
-                     ocl_dev_megakernel_program(rhs.ocl_dev_megakernel_program)
-               {
-                       /* copy can only happen in map insert, assert that */
-                       assert(mutex == NULL);
-               }
-
-               ~Slot()
-               {
-                       delete mutex;
-                       mutex = NULL;
-               }
-       };
-
-       /* key is combination of platform ID and device ID */
-       typedef pair<cl_platform_id, cl_device_id> PlatformDevicePair;
-
-       /* map of Slot objects */
-       typedef map<PlatformDevicePair, Slot> CacheMap;
-       CacheMap cache;
-
-       thread_mutex cache_lock;
-
-       /* lazy instantiate */
-       static OpenCLCache &global_instance()
-       {
-               static OpenCLCache instance;
-               return instance;
-       }
-
-       OpenCLCache()
-       {
-       }
-
-       ~OpenCLCache()
-       {
-               /* Intel OpenCL bug raises SIGABRT due to pure virtual call
-                * so this is disabled. It's not necessary to free objects
-                * at process exit anyway.
-                * http://software.intel.com/en-us/forums/topic/370083#comments */
-
-               //flush();
-       }
-
-       /* lookup something in the cache. If this returns NULL, slot_locker
-        * will be holding a lock for the cache. slot_locker should refer to a
-        * default constructed thread_scoped_lock */
-       template<typename T>
-       static T get_something(cl_platform_id platform,
-                              cl_device_id device,
-                              T Slot::*member,
-                              thread_scoped_lock& slot_locker)
-       {
-               assert(platform != NULL);
-
-               OpenCLCache& self = global_instance();
-
-               thread_scoped_lock cache_lock(self.cache_lock);
-
-               pair<CacheMap::iterator,bool> ins = self.cache.insert(
-                       CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
-
-               Slot &slot = ins.first->second;
-
-               /* create slot lock only while holding cache lock */
-               if(!slot.mutex)
-                       slot.mutex = new thread_mutex;
-
-               /* need to unlock cache before locking slot, to allow store to complete */
-               cache_lock.unlock();
-
-               /* lock the slot */
-               slot_locker = thread_scoped_lock(*slot.mutex);
-
-               /* If the thing isn't cached */
-               if(slot.*member == NULL) {
-                       /* return with the caller's lock holder holding the slot lock */
-                       return NULL;
-               }
-
-               /* the item was already cached, release the slot lock */
-               slot_locker.unlock();
-
-               return slot.*member;
-       }
-
-       /* store something in the cache. you MUST have tried to get the item before storing to it */
-       template<typename T>
-       static void store_something(cl_platform_id platform,
-                                   cl_device_id device,
-                                   T thing,
-                                   T Slot::*member,
-                                   thread_scoped_lock& slot_locker)
-       {
-               assert(platform != NULL);
-               assert(device != NULL);
-               assert(thing != NULL);
-
-               OpenCLCache &self = global_instance();
-
-               thread_scoped_lock cache_lock(self.cache_lock);
-               CacheMap::iterator i = self.cache.find(PlatformDevicePair(platform, device));
-               cache_lock.unlock();
-
-               Slot &slot = i->second;
-
-               /* sanity check */
-               assert(i != self.cache.end());
-               assert(slot.*member == NULL);
-
-               slot.*member = thing;
-
-               /* unlock the slot */
-               slot_locker.unlock();
-       }
-
-public:
-
-       enum ProgramName {
-               OCL_DEV_BASE_PROGRAM,
-               OCL_DEV_MEGAKERNEL_PROGRAM,
-       };
-
-       /* see get_something comment */
-       static cl_context get_context(cl_platform_id platform,
-                                     cl_device_id device,
-                                     thread_scoped_lock& slot_locker)
-       {
-               cl_context context = get_something<cl_context>(platform,
-                                                              device,
-                                                              &Slot::context,
-                                                              slot_locker);
-
-               if(!context)
-                       return NULL;
-
-               /* caller is going to release it when done with it, so retain it */
-               cl_int ciErr = clRetainContext(context);
-               assert(ciErr == CL_SUCCESS);
-               (void)ciErr;
-
-               return context;
-       }
-
-       /* see get_something comment */
-       static cl_program get_program(cl_platform_id platform,
-                                     cl_device_id device,
-                                     ProgramName program_name,
-                                     thread_scoped_lock& slot_locker)
-       {
-               cl_program program = NULL;
-
-               switch(program_name) {
-                       case OCL_DEV_BASE_PROGRAM:
-                               /* Get program related to OpenCLDeviceBase */
-                               program = get_something<cl_program>(platform,
-                                                                   device,
-                                                                   &Slot::ocl_dev_base_program,
-                                                                   slot_locker);
-                               break;
-                       case OCL_DEV_MEGAKERNEL_PROGRAM:
-                               /* Get program related to megakernel */
-                               program = get_something<cl_program>(platform,
-                                                                   device,
-                                                                   &Slot::ocl_dev_megakernel_program,
-                                                                   slot_locker);
-                               break;
-               default:
-                       assert(!"Invalid program name");
-               }
-
-               if(!program)
-                       return NULL;
-
-               /* caller is going to release it when done with it, so retain it */
-               cl_int ciErr = clRetainProgram(program);
-               assert(ciErr == CL_SUCCESS);
-               (void)ciErr;
-
-               return program;
-       }
-
-       /* see store_something comment */
-       static void store_context(cl_platform_id platform,
-                                 cl_device_id device,
-                                 cl_context context,
-                                 thread_scoped_lock& slot_locker)
-       {
-               store_something<cl_context>(platform,
-                                           device,
-                                           context,
-                                           &Slot::context,
-                                           slot_locker);
-
-               /* increment reference count in OpenCL.
-                * The caller is going to release the object when done with it. */
-               cl_int ciErr = clRetainContext(context);
-               assert(ciErr == CL_SUCCESS);
-               (void)ciErr;
-       }
-
-       /* see store_something comment */
-       static void store_program(cl_platform_id platform,
-                                 cl_device_id device,
-                                 cl_program program,
-                                 ProgramName program_name,
-                                 thread_scoped_lock& slot_locker)
-       {
-               switch(program_name) {
-                       case OCL_DEV_BASE_PROGRAM:
-                               store_something<cl_program>(platform,
-                                                           device,
-                                                           program,
-                                                           &Slot::ocl_dev_base_program,
-                                                           slot_locker);
-                               break;
-                       case OCL_DEV_MEGAKERNEL_PROGRAM:
-                               store_something<cl_program>(platform,
-                                                           device,
-                                                           program,
-                                                           &Slot::ocl_dev_megakernel_program,
-                                                           slot_locker);
-                               break;
-                       default:
-                               assert(!"Invalid program name\n");
-                               return;
-               }
-
-               /* Increment reference count in OpenCL.
-                * The caller is going to release the object when done with it.
-                */
-               cl_int ciErr = clRetainProgram(program);
-               assert(ciErr == CL_SUCCESS);
-               (void)ciErr;
-       }
-
-       /* Discard all cached contexts and programs.  */
-       static void flush()
-       {
-               OpenCLCache &self = global_instance();
-               thread_scoped_lock cache_lock(self.cache_lock);
-
-               foreach(CacheMap::value_type &item, self.cache) {
-                       if(item.second.ocl_dev_base_program != NULL)
-                               clReleaseProgram(item.second.ocl_dev_base_program);
-                       if(item.second.ocl_dev_megakernel_program != NULL)
-                               clReleaseProgram(item.second.ocl_dev_megakernel_program);
-                       if(item.second.context != NULL)
-                               clReleaseContext(item.second.context);
-               }
-
-               self.cache.clear();
-       }
-};
-
-class OpenCLDeviceBase : public Device
-{
-public:
-       DedicatedTaskPool task_pool;
-       cl_context cxContext;
-       cl_command_queue cqCommandQueue;
-       cl_platform_id cpPlatform;
-       cl_device_id cdDevice;
-       cl_program cpProgram;
-       cl_kernel ckFilmConvertByteKernel;
-       cl_kernel ckFilmConvertHalfFloatKernel;
-       cl_kernel ckShaderKernel;
-       cl_kernel ckBakeKernel;
-       cl_int ciErr;
-
-       typedef map<string, device_vector<uchar>*> ConstMemMap;
-       typedef map<string, device_ptr> MemMap;
-
-       ConstMemMap const_mem_map;
-       MemMap mem_map;
-       device_ptr null_mem;
-
-       bool device_initialized;
-       string platform_name;
-
-       bool opencl_error(cl_int err)
-       {
-               if(err != CL_SUCCESS) {
-                       string message = string_printf("OpenCL error (%d): %s", err, clewErrorString(err));
-                       if(error_msg == "")
-                               error_msg = message;
-                       fprintf(stderr, "%s\n", message.c_str());
-                       return true;
-               }
-
-               return false;
-       }
-
-       void opencl_error(const string& message)
-       {
-               if(error_msg == "")
-                       error_msg = message;
-               fprintf(stderr, "%s\n", message.c_str());
-       }
-
-#define opencl_assert(stmt) \
-       { \
-               cl_int err = stmt; \
-               \
-               if(err != CL_SUCCESS) { \
-                       string message = string_printf("OpenCL error: %s in %s", clewErrorString(err), #stmt); \
-                       if(error_msg == "") \
-                               error_msg = message; \
-                       fprintf(stderr, "%s\n", message.c_str()); \
-               } \
-       } (void)0
-
-       void opencl_assert_err(cl_int err, const char* where)
-       {
-               if(err != CL_SUCCESS) {
-                       string message = string_printf("OpenCL error (%d): %s in %s", err, clewErrorString(err), where);
-                       if(error_msg == "")
-                               error_msg = message;
-                       fprintf(stderr, "%s\n", message.c_str());
-#ifndef NDEBUG
-                       abort();
-#endif
-               }
-       }
-
-       OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
-       : Device(info, stats, background_)
-       {
-               cpPlatform = NULL;
-               cdDevice = NULL;
-               cxContext = NULL;
-               cqCommandQueue = NULL;
-               cpProgram = NULL;
-               ckFilmConvertByteKernel = NULL;
-               ckFilmConvertHalfFloatKernel = NULL;
-               ckShaderKernel = NULL;
-               ckBakeKernel = NULL;
-               null_mem = 0;
-               device_initialized = false;
-
-               vector<OpenCLPlatformDevice> usable_devices;
-               opencl_get_usable_devices(&usable_devices);
-               if(usable_devices.size() == 0) {
-                       opencl_error("OpenCL: no devices found.");
-                       return;
-               }
-               assert(info.num < usable_devices.size());
-               OpenCLPlatformDevice& platform_device = usable_devices[info.num];
-               cpPlatform = platform_device.platform_id;
-               cdDevice = platform_device.device_id;
-               platform_name = platform_device.platform_name;
-               VLOG(2) << "Creating new Cycles device for OpenCL platform "
-                       << platform_name << ", device "
-                       << platform_device.device_name << ".";
-
-               {
-                       /* try to use cached context */
-                       thread_scoped_lock cache_locker;
-                       cxContext = OpenCLCache::get_context(cpPlatform, cdDevice, cache_locker);
-
-                       if(cxContext == NULL) {
-                               /* create context properties array to specify platform */
-                               const cl_context_properties context_props[] = {
-                                       CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform,
-                                       0, 0
-                               };
-
-                               /* create context */
-                               cxContext = clCreateContext(context_props, 1, &cdDevice,
-                                       context_notify_callback, cdDevice, &ciErr);
-
-                               if(opencl_error(ciErr)) {
-                                       opencl_error("OpenCL: clCreateContext failed");
-                                       return;
-                               }
-
-                               /* cache it */
-                               OpenCLCache::store_context(cpPlatform, cdDevice, cxContext, cache_locker);
-                       }
-               }
-
-               cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
-               if(opencl_error(ciErr))
-                       return;
-
-               null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
-               if(opencl_error(ciErr))
-                       return;
-
-               fprintf(stderr, "Device init success\n");
-               device_initialized = true;
-       }
-
-       static void CL_CALLBACK context_notify_callback(const char *err_info,
-               const void * /*private_info*/, size_t /*cb*/, void *user_data)
-       {
-               char name[256];
-               clGetDeviceInfo((cl_device_id)user_data, CL_DEVICE_NAME, sizeof(name), &name, NULL);
-
-               fprintf(stderr, "OpenCL error (%s): %s\n", name, err_info);
-       }
-
-       bool opencl_version_check()
-       {
-               string error;
-               if(!opencl_platform_version_check(cpPlatform, &error)) {
-                       opencl_error(error);
-                       return false;
-               }
-               if(!opencl_device_version_check(cdDevice, &error)) {
-                       opencl_error(error);
-                       return false;
-               }
-               return true;
-       }
-
-       bool load_binary(const string& /*kernel_path*/,
-                        const string& clbin,
-                        const string& custom_kernel_build_options,
-                        cl_program *program,
-                        const string *debug_src = NULL)
-       {
-               /* read binary into memory */
-               vector<uint8_t> binary;
-
-               if(!path_read_binary(clbin, binary)) {
-                       opencl_error(string_printf("OpenCL failed to read cached binary %s.", clbin.c_str()));
-                       return false;
-               }
-
-               /* create program */
-               cl_int status;
-               size_t size = binary.size();
-               const uint8_t *bytes = &binary[0];
-
-               *program = clCreateProgramWithBinary(cxContext, 1, &cdDevice,
-                       &size, &bytes, &status, &ciErr);
-
-               if(opencl_error(status) || opencl_error(ciErr)) {
-                       opencl_error(string_printf("OpenCL failed create program from cached binary %s.", clbin.c_str()));
-                       return false;
-               }
-
-               if(!build_kernel(program, custom_kernel_build_options, debug_src))
-                       return false;
-
-               return true;
-       }
-
-       bool save_binary(cl_program *program, const string& clbin)
-       {
-               size_t size = 0;
-               clGetProgramInfo(*program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
-
-               if(!size)
-                       return false;
-
-               vector<uint8_t> binary(size);
-               uint8_t *bytes = &binary[0];
-
-               clGetProgramInfo(*program, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
-
-               if(!path_write_binary(clbin, binary)) {
-                       opencl_error(string_printf("OpenCL failed to write cached binary %s.", clbin.c_str()));
-                       return false;
-               }
-
-               return true;
-       }
-
-       bool build_kernel(cl_program *kernel_program,
-                         const string& custom_kernel_build_options,
-                         const string *debug_src = NULL)
-       {
-               string build_options;
-               build_options = kernel_build_options(debug_src) + custom_kernel_build_options;
-
-               ciErr = clBuildProgram(*kernel_program, 0, NULL, build_options.c_str(), NULL, NULL);
-
-               /* show warnings even if build is successful */
-               size_t ret_val_size = 0;
-
-               clGetProgramBuildInfo(*kernel_program, cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
-
-               if(ret_val_size > 1) {
-                       vector<char> build_log(ret_val_size + 1);
-                       clGetProgramBuildInfo(*kernel_program, cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, &build_log[0], NULL);
-
-                       build_log[ret_val_size] = '\0';
-                       /* Skip meaningless empty output from the NVidia compiler. */
-                       if(!(ret_val_size == 2 && build_log[0] == '\n')) {
-                               fprintf(stderr, "OpenCL kernel build output:\n");
-                               fprintf(stderr, "%s\n", &build_log[0]);
-                       }
-               }
-
-               if(ciErr != CL_SUCCESS) {
-                       opencl_error("OpenCL build failed: errors in console");
-                       fprintf(stderr, "Build error: %s\n", clewErrorString(ciErr));
-                       return false;
-               }
-
-               return true;
-       }
-
-       bool compile_kernel(const string& kernel_name,
-                           const string& kernel_path,
-                           const string& source,
-                           const string& custom_kernel_build_options,
-                           cl_program *kernel_program,
-                           const string *debug_src = NULL)
-       {
-               /* We compile kernels consisting of many files. unfortunately OpenCL
-                * kernel caches do not seem to recognize changes in included files.
-                * so we force recompile on changes by adding the md5 hash of all files.
-                */
-               string inlined_source = path_source_replace_includes(source,
-                                                                    kernel_path);
-
-               if(debug_src) {
-                       path_write_text(*debug_src, inlined_source);
-               }
-
-               size_t source_len = inlined_source.size();
-               const char *source_str = inlined_source.c_str();
-
-               *kernel_program = clCreateProgramWithSource(cxContext,
-                                                           1,
-                                                           &source_str,
-                                                           &source_len,
-                                                           &ciErr);
-
-               if(opencl_error(ciErr)) {
-                       return false;
-               }
-
-               double starttime = time_dt();
-               printf("Compiling %s OpenCL kernel ...\n", kernel_name.c_str());
-               /* TODO(sergey): Report which kernel is being compiled
-                * as well (megakernel or which of split kernels etc..).
-                */
-               printf("Build flags: %s\n", custom_kernel_build_options.c_str());
-
-               if(!build_kernel(kernel_program, custom_kernel_build_options, debug_src))
-                       return false;
-
-               printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
-
-               return true;
-       }
-
-       string device_md5_hash(string kernel_custom_build_options = "")
-       {
-               MD5Hash md5;
-               char version[256], driver[256], name[256], vendor[256];
-
-               clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
-               clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
-               clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
-               clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
-
-               md5.append((uint8_t*)vendor, strlen(vendor));
-               md5.append((uint8_t*)version, strlen(version));
-               md5.append((uint8_t*)name, strlen(name));
-               md5.append((uint8_t*)driver, strlen(driver));
-
-               string options = kernel_build_options();
-               options += kernel_custom_build_options;
-               md5.append((uint8_t*)options.c_str(), options.size());
-
-               return md5.get_hex();
-       }
-
-       bool load_kernels(const DeviceRequestedFeatures& requested_features)
-       {
-               /* Verify if device was initialized. */
-               if(!device_initialized) {
-                       fprintf(stderr, "OpenCL: failed to initialize device.\n");
-                       return false;
-               }
-
-               /* Try to use cached kernel. */
-               thread_scoped_lock cache_locker;
-               cpProgram = load_cached_kernel(requested_features,
-                                              OpenCLCache::OCL_DEV_BASE_PROGRAM,
-                                              cache_locker);
-
-               if(!cpProgram) {
-                       VLOG(2) << "No cached OpenCL kernel.";
-
-                       /* Verify we have right opencl version. */
-                       if(!opencl_version_check())
-                               return false;
-
-                       string build_flags = build_options_for_base_program(requested_features);
-
-                       /* Calculate md5 hashes to detect changes. */
-                       string kernel_path = path_get("kernel");
-                       string kernel_md5 = path_files_md5_hash(kernel_path);
-                       string device_md5 = device_md5_hash(build_flags);
-
-                       /* Path to cached binary.
-                        *
-                        * TODO(sergey): Seems we could de-duplicate all this string_printf()
-                        * calls with some utility function which will give file name for a
-                        * given hashes..
-                        */
-                       string clbin = string_printf("cycles_kernel_%s_%s.clbin",
-                                                    device_md5.c_str(),
-                                                    kernel_md5.c_str());
-                       clbin = path_cache_get(path_join("kernels", clbin));
-
-                       /* path to preprocessed source for debugging */
-                       string clsrc, *debug_src = NULL;
-
-                       if(opencl_kernel_use_debug()) {
-                               clsrc = string_printf("cycles_kernel_%s_%s.cl",
-                                                     device_md5.c_str(),
-                                                     kernel_md5.c_str());
-                               clsrc = path_cache_get(path_join("kernels", clsrc));
-                               debug_src = &clsrc;
-                       }
-
-                       /* If binary kernel exists already, try use it. */
-                       if(path_exists(clbin) && load_binary(kernel_path,
-                                                            clbin,
-                                                            build_flags,
-                                                            &cpProgram))
-                       {
-                               /* Kernel loaded from binary, nothing to do. */
-                               VLOG(2) << "Loaded kernel from " << clbin << ".";
-                       }
-                       else {
-                               VLOG(2) << "Kernel file " << clbin << " either doesn't exist or failed to be loaded by driver.";
-                               string init_kernel_source = "#include \"kernels/opencl/kernel.cl\" // " + kernel_md5 + "\n";
-
-                               /* If does not exist or loading binary failed, compile kernel. */
-                               if(!compile_kernel("base_kernel",
-                                                  kernel_path,
-                                                  init_kernel_source,
-                                                  build_flags,
-                                                  &cpProgram,
-                                                  debug_src))
-                               {
-                                       return false;
-                               }
-
-                               /* Save binary for reuse. */
-                               if(!save_binary(&cpProgram, clbin)) {
-                                       return false;
-                               }
-                       }
-
-                       /* Cache the program. */
-                       store_cached_kernel(cpPlatform,
-                                           cdDevice,
-                                           cpProgram,
-                                           OpenCLCache::OCL_DEV_BASE_PROGRAM,
-                                           cache_locker);
-               }
-               else {
-                       VLOG(2) << "Found cached OpenCL kernel.";
-               }
-
-               /* Find kernels. */
-#define FIND_KERNEL(kernel_var, kernel_name) \
-               do { \
-                       kernel_var = clCreateKernel(cpProgram, "kernel_ocl_" kernel_name, &ciErr); \
-                       if(opencl_error(ciErr)) \
-                               return false; \
-               } while(0)
-
-               FIND_KERNEL(ckFilmConvertByteKernel, "convert_to_byte");
-               FIND_KERNEL(ckFilmConvertHalfFloatKernel, "convert_to_half_float");
-               FIND_KERNEL(ckShaderKernel, "shader");
-               FIND_KERNEL(ckBakeKernel, "bake");
-
-#undef FIND_KERNEL
-               return true;
-       }
-
-       ~OpenCLDeviceBase()
-       {
-               task_pool.stop();
-
-               if(null_mem)
-                       clReleaseMemObject(CL_MEM_PTR(null_mem));
-
-               ConstMemMap::iterator mt;
-               for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
-                       mem_free(*(mt->second));
-                       delete mt->second;
-               }
-
-               if(ckFilmConvertByteKernel)
-                       clReleaseKernel(ckFilmConvertByteKernel);
-               if(ckFilmConvertHalfFloatKernel)
-                       clReleaseKernel(ckFilmConvertHalfFloatKernel);
-               if(ckShaderKernel)
-                       clReleaseKernel(ckShaderKernel);
-               if(ckBakeKernel)
-                       clReleaseKernel(ckBakeKernel);
-               if(cpProgram)
-                       clReleaseProgram(cpProgram);
-               if(cqCommandQueue)
-                       clReleaseCommandQueue(cqCommandQueue);
-               if(cxContext)
-                       clReleaseContext(cxContext);
-       }
-
-       void mem_alloc(device_memory& mem, MemoryType type)
-       {
-               size_t size = mem.memory_size();
-
-               cl_mem_flags mem_flag;
-               void *mem_ptr = NULL;
-
-               if(type == MEM_READ_ONLY)
-                       mem_flag = CL_MEM_READ_ONLY;
-               else if(type == MEM_WRITE_ONLY)
-                       mem_flag = CL_MEM_WRITE_ONLY;
-               else
-                       mem_flag = CL_MEM_READ_WRITE;
-
-               /* Zero-size allocation might be invoked by render, but not really
-                * supported by OpenCL. Using NULL as device pointer also doesn't really
-                * work for some reason, so for the time being we'll use special case
-                * will null_mem buffer.
-                */
-               if(size != 0) {
-                       mem.device_pointer = (device_ptr)clCreateBuffer(cxContext,
-                                                                       mem_flag,
-                                                                       size,
-                                                                       mem_ptr,
-                                                                       &ciErr);
-                       opencl_assert_err(ciErr, "clCreateBuffer");
-               }
-               else {
-                       mem.device_pointer = null_mem;
-               }
-
-               stats.mem_alloc(size);
-               mem.device_size = size;
-       }
-
-       void mem_copy_to(device_memory& mem)
-       {
-               /* this is blocking */
-               size_t size = mem.memory_size();
-               if(size != 0) {
-                       opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
-                                                          CL_MEM_PTR(mem.device_pointer),
-                                                          CL_TRUE,
-                                                          0,
-                                                          size,
-                                                          (void*)mem.data_pointer,
-                                                          0,
-                                                          NULL, NULL));
-               }
-       }
-
-       void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
-       {
-               size_t offset = elem*y*w;
-               size_t size = elem*w*h;
-               assert(size != 0);
-               opencl_assert(clEnqueueReadBuffer(cqCommandQueue,
-                                                 CL_MEM_PTR(mem.device_pointer),
-                                                 CL_TRUE,
-                                                 offset,
-                                                 size,
-                                                 (uchar*)mem.data_pointer + offset,
-                                                 0,
-                                                 NULL, NULL));
-       }
-
-       void mem_zero(device_memory& mem)
-       {
-               if(mem.device_pointer) {
-                       memset((void*)mem.data_pointer, 0, mem.memory_size());
-                       mem_copy_to(mem);
-               }
-       }
-
-       void mem_free(device_memory& mem)
-       {
-               if(mem.device_pointer) {
-                       if(mem.device_pointer != null_mem) {
-                               opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
-                       }
-                       mem.device_pointer = 0;
-
-                       stats.mem_free(mem.device_size);
-                       mem.device_size = 0;
-               }
-       }
-
-       void const_copy_to(const char *name, void *host, size_t size)
-       {
-               ConstMemMap::iterator i = const_mem_map.find(name);
-
-               if(i == const_mem_map.end()) {
-                       device_vector<uchar> *data = new device_vector<uchar>();
-                       data->copy((uchar*)host, size);
-
-                       mem_alloc(*data, MEM_READ_ONLY);
-                       i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first;
-               }
-               else {
-                       device_vector<uchar> *data = i->second;
-                       data->copy((uchar*)host, size);
-               }
-
-               mem_copy_to(*i->second);
-       }
-
-       void tex_alloc(const char *name,
-                      device_memory& mem,
-                      InterpolationType /*interpolation*/,
-                      ExtensionType /*extension*/)
-       {
-               VLOG(1) << "Texture allocate: " << name << ", "
-                       << string_human_readable_number(mem.memory_size()) << " bytes. ("
-                       << string_human_readable_size(mem.memory_size()) << ")";
-               mem_alloc(mem, MEM_READ_ONLY);
-               mem_copy_to(mem);
-               assert(mem_map.find(name) == mem_map.end());
-               mem_map.insert(MemMap::value_type(name, mem.device_pointer));
-       }
-
-       void tex_free(device_memory& mem)
-       {
-               if(mem.device_pointer) {
-                       foreach(const MemMap::value_type& value, mem_map) {
-                               if(value.second == mem.device_pointer) {
-                                       mem_map.erase(value.first);
-                                       break;
-                               }
-                       }
-
-                       mem_free(mem);
-               }
-       }
-
-       size_t global_size_round_up(int group_size, int global_size)
-       {
-               int r = global_size % group_size;
-               return global_size + ((r == 0)? 0: group_size - r);
-       }
-
-       void enqueue_kernel(cl_kernel kernel, size_t w, size_t h)
-       {
-               size_t workgroup_size, max_work_items[3];
-
-               clGetKernelWorkGroupInfo(kernel, cdDevice,
-                       CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
-               clGetDeviceInfo(cdDevice,
-                       CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, max_work_items, NULL);
-
-               /* Try to divide evenly over 2 dimensions. */
-               size_t sqrt_workgroup_size = max((size_t)sqrt((double)workgroup_size), 1);
-               size_t local_size[2] = {sqrt_workgroup_size, sqrt_workgroup_size};
-
-               /* Some implementations have max size 1 on 2nd dimension. */
-               if(local_size[1] > max_work_items[1]) {
-                       local_size[0] = workgroup_size/max_work_items[1];
-                       local_size[1] = max_work_items[1];
-               }
-
-               size_t global_size[2] = {global_size_round_up(local_size[0], w),
-                                        global_size_round_up(local_size[1], h)};
-
-               /* Vertical size of 1 is coming from bake/shade kernels where we should
-                * not round anything up because otherwise we'll either be doing too
-                * much work per pixel (if we don't check global ID on Y axis) or will
-                * be checking for global ID to always have Y of 0.
-                */
-               if (h == 1) {
-                       global_size[h] = 1;
-               }
-
-               /* run kernel */
-               opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));
-               opencl_assert(clFlush(cqCommandQueue));
-       }
-
-       void set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name)
-       {
-               cl_mem ptr;
-
-               MemMap::iterator i = mem_map.find(name);
-               if(i != mem_map.end()) {
-                       ptr = CL_MEM_PTR(i->second);
-               }
-               else {
-                       /* work around NULL not working, even though the spec says otherwise */
-                       ptr = CL_MEM_PTR(null_mem);
-               }
-               
-               opencl_assert(clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr));
-       }
-
-       void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
-       {
-               /* cast arguments to cl types */
-               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
-               cl_mem d_rgba = (rgba_byte)? CL_MEM_PTR(rgba_byte): CL_MEM_PTR(rgba_half);
-               cl_mem d_buffer = CL_MEM_PTR(buffer);
-               cl_int d_x = task.x;
-               cl_int d_y = task.y;
-               cl_int d_w = task.w;
-               cl_int d_h = task.h;
-               cl_float d_sample_scale = 1.0f/(task.sample + 1);
-               cl_int d_offset = task.offset;
-               cl_int d_stride = task.stride;
-
-
-               cl_kernel ckFilmConvertKernel = (rgba_byte)? ckFilmConvertByteKernel: ckFilmConvertHalfFloatKernel;
-
-               cl_uint start_arg_index =
-                       kernel_set_args(ckFilmConvertKernel,
-                                       0,
-                                       d_data,
-                                       d_rgba,
-                                       d_buffer);
-
-#define KERNEL_TEX(type, ttype, name) \
-       set_kernel_arg_mem(ckFilmConvertKernel, &start_arg_index, #name);
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-
-               start_arg_index += kernel_set_args(ckFilmConvertKernel,
-                                                  start_arg_index,
-                                                  d_sample_scale,
-                                                  d_x,
-                                                  d_y,
-                                                  d_w,
-                                                  d_h,
-                                                  d_offset,
-                                                  d_stride);
-
-               enqueue_kernel(ckFilmConvertKernel, d_w, d_h);
-       }
-
-       void shader(DeviceTask& task)
-       {
-               /* cast arguments to cl types */
-               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
-               cl_mem d_input = CL_MEM_PTR(task.shader_input);
-               cl_mem d_output = CL_MEM_PTR(task.shader_output);
-               cl_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma);
-               cl_int d_shader_eval_type = task.shader_eval_type;
-               cl_int d_shader_filter = task.shader_filter;
-               cl_int d_shader_x = task.shader_x;
-               cl_int d_shader_w = task.shader_w;
-               cl_int d_offset = task.offset;
-
-               cl_kernel kernel;
-
-               if(task.shader_eval_type >= SHADER_EVAL_BAKE)
-                       kernel = ckBakeKernel;
-               else
-                       kernel = ckShaderKernel;
-
-               cl_uint start_arg_index =
-                       kernel_set_args(kernel,
-                                       0,
-                                       d_data,
-                                       d_input,
-                                       d_output);
-
-               if(task.shader_eval_type < SHADER_EVAL_BAKE) {
-                       start_arg_index += kernel_set_args(kernel,
-                                                          start_arg_index,
-                                                          d_output_luma);
-               }
-
-#define KERNEL_TEX(type, ttype, name) \
-               set_kernel_arg_mem(kernel, &start_arg_index, #name);
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-
-               start_arg_index += kernel_set_args(kernel,
-                                                  start_arg_index,
-                                                  d_shader_eval_type);
-               if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
-                       start_arg_index += kernel_set_args(kernel,
-                                                          start_arg_index,
-                                                          d_shader_filter);
-               }
-               start_arg_index += kernel_set_args(kernel,
-                                                  start_arg_index,
-                                                  d_shader_x,
-                                                  d_shader_w,
-                                                  d_offset);
-
-               for(int sample = 0; sample < task.num_samples; sample++) {
-
-                       if(task.get_cancel())
-                               break;
-
-                       kernel_set_args(kernel, start_arg_index, sample);
-
-                       enqueue_kernel(kernel, task.shader_w, 1);
-
-                       clFinish(cqCommandQueue);
-
-                       task.update_progress(NULL);
-               }
-       }
-
-       class OpenCLDeviceTask : public DeviceTask {
-       public:
-               OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
-               : DeviceTask(task)
-               {
-                       run = function_bind(&OpenCLDeviceBase::thread_run,
-                                           device,
-                                           this);
-               }
-       };
-
-       int get_split_task_count(DeviceTask& /*task*/)
-       {
-               return 1;
-       }
-
-       void task_add(DeviceTask& task)
-       {
-               task_pool.push(new OpenCLDeviceTask(this, task));
-       }
-
-       void task_wait()
-       {
-               task_pool.wait();
-       }
-
-       void task_cancel()
-       {
-               task_pool.cancel();
-       }
-
-       virtual void thread_run(DeviceTask * /*task*/) = 0;
-
-protected:
-       string kernel_build_options(const string *debug_src = NULL)
-       {
-               string build_options = "-cl-fast-relaxed-math ";
-
-               if(platform_name == "NVIDIA CUDA") {
-                       build_options += "-D__KERNEL_OPENCL_NVIDIA__ "
-                                        "-cl-nv-maxrregcount=32 "
-                                        "-cl-nv-verbose ";
-
-                       uint compute_capability_major, compute_capability_minor;
-                       clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,
-                                       sizeof(cl_uint), &compute_capability_major, NULL);
-                       clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV,
-                                       sizeof(cl_uint), &compute_capability_minor, NULL);
-
-                       build_options += string_printf("-D__COMPUTE_CAPABILITY__=%u ",
-                                                      compute_capability_major * 100 +
-                                                      compute_capability_minor * 10);
-               }
-
-               else if(platform_name == "Apple")
-                       build_options += "-D__KERNEL_OPENCL_APPLE__ ";
-
-               else if(platform_name == "AMD Accelerated Parallel Processing")
-                       build_options += "-D__KERNEL_OPENCL_AMD__ ";
-
-               else if(platform_name == "Intel(R) OpenCL") {
-                       build_options += "-D__KERNEL_OPENCL_INTEL_CPU__ ";
-
-                       /* Options for gdb source level kernel debugging.
-                        * this segfaults on linux currently.
-                        */
-                       if(opencl_kernel_use_debug() && debug_src)
-                               build_options += "-g -s \"" + *debug_src + "\" ";
-               }
-
-               if(opencl_kernel_use_debug())
-                       build_options += "-D__KERNEL_OPENCL_DEBUG__ ";
-
-#ifdef WITH_CYCLES_DEBUG
-               build_options += "-D__KERNEL_DEBUG__ ";
-#endif
-
-               return build_options;
-       }
-
-       class ArgumentWrapper {
-       public:
-               ArgumentWrapper() : size(0), pointer(NULL) {}
-               template <typename T>
-               ArgumentWrapper(T& argument) : size(sizeof(argument)),
-                                              pointer(&argument) { }
-               ArgumentWrapper(int argument) : size(sizeof(int)),
-                                               int_value(argument),
-                                               pointer(&int_value) { }
-               ArgumentWrapper(float argument) : size(sizeof(float)),
-                                                 float_value(argument),
-                                                 pointer(&float_value) { }
-               size_t size;
-               int int_value;
-               float float_value;
-               void *pointer;
-       };
-
-       /* TODO(sergey): In the future we can use variadic templates, once
-        * C++0x is allowed. Should allow to clean this up a bit.
-        */
-       int kernel_set_args(cl_kernel kernel,
-                           int start_argument_index,
-                           const ArgumentWrapper& arg1 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg2 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg3 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg4 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg5 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg6 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg7 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg8 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg9 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg10 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg11 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg12 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg13 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg14 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg15 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg16 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg17 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg18 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg19 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg20 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg21 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg22 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg23 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg24 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg25 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg26 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg27 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg28 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg29 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg30 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg31 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg32 = ArgumentWrapper(),
-                           const ArgumentWrapper& arg33 = ArgumentWrapper())
-       {
-               int current_arg_index = 0;
-#define FAKE_VARARG_HANDLE_ARG(arg) \
-               do { \
-                       if(arg.pointer != NULL) { \
-                               opencl_assert(clSetKernelArg( \
-                                       kernel, \
-                                       start_argument_index + current_arg_index, \
-                                       arg.size, arg.pointer)); \
-                               ++current_arg_index; \
-                       } \
-                       else { \
-                               return current_arg_index; \
-                       } \
-               } while(false)
-               FAKE_VARARG_HANDLE_ARG(arg1);
-               FAKE_VARARG_HANDLE_ARG(arg2);
-               FAKE_VARARG_HANDLE_ARG(arg3);
-               FAKE_VARARG_HANDLE_ARG(arg4);
-               FAKE_VARARG_HANDLE_ARG(arg5);
-               FAKE_VARARG_HANDLE_ARG(arg6);
-               FAKE_VARARG_HANDLE_ARG(arg7);
-               FAKE_VARARG_HANDLE_ARG(arg8);
-               FAKE_VARARG_HANDLE_ARG(arg9);
-               FAKE_VARARG_HANDLE_ARG(arg10);
-               FAKE_VARARG_HANDLE_ARG(arg11);
-               FAKE_VARARG_HANDLE_ARG(arg12);
-               FAKE_VARARG_HANDLE_ARG(arg13);
-               FAKE_VARARG_HANDLE_ARG(arg14);
-               FAKE_VARARG_HANDLE_ARG(arg15);
-               FAKE_VARARG_HANDLE_ARG(arg16);
-               FAKE_VARARG_HANDLE_ARG(arg17);
-               FAKE_VARARG_HANDLE_ARG(arg18);
-               FAKE_VARARG_HANDLE_ARG(arg19);
-               FAKE_VARARG_HANDLE_ARG(arg20);
-               FAKE_VARARG_HANDLE_ARG(arg21);
-               FAKE_VARARG_HANDLE_ARG(arg22);
-               FAKE_VARARG_HANDLE_ARG(arg23);
-               FAKE_VARARG_HANDLE_ARG(arg24);
-               FAKE_VARARG_HANDLE_ARG(arg25);
-               FAKE_VARARG_HANDLE_ARG(arg26);
-               FAKE_VARARG_HANDLE_ARG(arg27);
-               FAKE_VARARG_HANDLE_ARG(arg28);
-               FAKE_VARARG_HANDLE_ARG(arg29);
-               FAKE_VARARG_HANDLE_ARG(arg30);
-               FAKE_VARARG_HANDLE_ARG(arg31);
-               FAKE_VARARG_HANDLE_ARG(arg32);
-               FAKE_VARARG_HANDLE_ARG(arg33);
-#undef FAKE_VARARG_HANDLE_ARG
-               return current_arg_index;
-       }
-
-       inline void release_kernel_safe(cl_kernel kernel)
-       {
-               if(kernel) {
-                       clReleaseKernel(kernel);
-               }
-       }
-
-       inline void release_mem_object_safe(cl_mem mem)
-       {
-               if(mem != NULL) {
-                       clReleaseMemObject(mem);
-               }
-       }
-
-       inline void release_program_safe(cl_program program)
-       {
-               if(program) {
-                       clReleaseProgram(program);
-               }
-       }
-
-       /* ** Those guys are for workign around some compiler-specific bugs ** */
-
-       virtual cl_program load_cached_kernel(
-               const DeviceRequestedFeatures& /*requested_features*/,
-               OpenCLCache::ProgramName program_name,
-               thread_scoped_lock& cache_locker)
-       {
-               return OpenCLCache::get_program(cpPlatform,
-                                               cdDevice,
-                                               program_name,
-                                               cache_locker);
-       }
-
-       virtual void store_cached_kernel(cl_platform_id platform,
-                                        cl_device_id device,
-                                        cl_program program,
-                                        OpenCLCache::ProgramName program_name,
-                                        thread_scoped_lock& cache_locker)
-       {
-               OpenCLCache::store_program(platform,
-                                          device,
-                                          program,
-                                          program_name,
-                                          cache_locker);
-       }
-
-       virtual string build_options_for_base_program(
-               const DeviceRequestedFeatures& /*requested_features*/)
-       {
-               /* TODO(sergey): By default we compile all features, meaning
-                * mega kernel is not getting feature-based optimizations.
-                *
-                * Ideally we need always compile kernel with as less features
-                * enabled as possible to keep performance at it's max.
-                */
-               return "";
-       }
-};
-
-class OpenCLDeviceMegaKernel : public OpenCLDeviceBase
-{
-public:
-       cl_kernel ckPathTraceKernel;
-       cl_program path_trace_program;
-
-       OpenCLDeviceMegaKernel(DeviceInfo& info, Stats &stats, bool background_)
-       : OpenCLDeviceBase(info, stats, background_)
-       {
-               ckPathTraceKernel = NULL;
-               path_trace_program = NULL;
-       }
-
-       bool load_kernels(const DeviceRequestedFeatures& requested_features)
-       {
-               /* Get Shader, bake and film convert kernels.
-                * It'll also do verification of OpenCL actually initialized.
-                */
-               if(!OpenCLDeviceBase::load_kernels(requested_features)) {
-                       return false;
-               }
-
-               /* Try to use cached kernel. */
-               thread_scoped_lock cache_locker;
-               path_trace_program = OpenCLCache::get_program(cpPlatform,
-                                                             cdDevice,
-                                                             OpenCLCache::OCL_DEV_MEGAKERNEL_PROGRAM,
-                                                             cache_locker);
-
-               if(!path_trace_program) {
-                       /* Verify we have right opencl version. */
-                       if(!opencl_version_check())
-                               return false;
-
-                       /* Calculate md5 hash to detect changes. */
-                       string kernel_path = path_get("kernel");
-                       string kernel_md5 = path_files_md5_hash(kernel_path);
-                       string custom_kernel_build_options = "-D__COMPILE_ONLY_MEGAKERNEL__ ";
-                       string device_md5 = device_md5_hash(custom_kernel_build_options);
-
-                       /* Path to cached binary. */
-                       string clbin = string_printf("cycles_kernel_%s_%s.clbin",
-                                                    device_md5.c_str(),
-                                                    kernel_md5.c_str());
-                       clbin = path_cache_get(path_join("kernels", clbin));
-
-                       /* Path to preprocessed source for debugging. */
-                       string clsrc, *debug_src = NULL;
-                       if(opencl_kernel_use_debug()) {
-                               clsrc = string_printf("cycles_kernel_%s_%s.cl",
-                                                     device_md5.c_str(),
-                                                     kernel_md5.c_str());
-                               clsrc = path_cache_get(path_join("kernels", clsrc));
-                               debug_src = &clsrc;
-                       }
-
-                       /* If exists already, try use it. */
-                       if(path_exists(clbin) && load_binary(kernel_path,
-                                                            clbin,
-                                                            custom_kernel_build_options,
-                                                            &path_trace_program,
-                                                            debug_src))
-                       {
-                               /* Kernel loaded from binary, nothing to do. */
-                       }
-                       else {
-                               string init_kernel_source = "#include \"kernels/opencl/kernel.cl\" // " +
-                                                           kernel_md5 + "\n";
-                               /* If does not exist or loading binary failed, compile kernel. */
-                               if(!compile_kernel("mega_kernel",
-                                                  kernel_path,
-                                                  init_kernel_source,
-                                                  custom_kernel_build_options,
-                                                  &path_trace_program,
-                                                  debug_src))
-                               {
-                                       return false;
-                               }
-                               /* Save binary for reuse. */
-                               if(!save_binary(&path_trace_program, clbin)) {
-                                       return false;
-                               }
-                       }
-                       /* Cache the program. */
-                       OpenCLCache::store_program(cpPlatform,
-                                                  cdDevice,
-                                                  path_trace_program,
-                                                  OpenCLCache::OCL_DEV_MEGAKERNEL_PROGRAM,
-                                                  cache_locker);
-               }
-
-               /* Find kernels. */
-               ckPathTraceKernel = clCreateKernel(path_trace_program,
-                                                  "kernel_ocl_path_trace",
-                                                  &ciErr);
-               if(opencl_error(ciErr))
-                       return false;
-               return true;
-       }
-
-       ~OpenCLDeviceMegaKernel()
-       {
-               task_pool.stop();
-               release_kernel_safe(ckPathTraceKernel);
-               release_program_safe(path_trace_program);
-       }
-
-       void path_trace(RenderTile& rtile, int sample)
-       {
-               /* Cast arguments to cl types. */
-               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
-               cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
-               cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
-               cl_int d_x = rtile.x;
-               cl_int d_y = rtile.y;
-               cl_int d_w = rtile.w;
-               cl_int d_h = rtile.h;
-               cl_int d_offset = rtile.offset;
-               cl_int d_stride = rtile.stride;
-
-               /* Sample arguments. */
-               cl_int d_sample = sample;
-
-               cl_uint start_arg_index =
-                       kernel_set_args(ckPathTraceKernel,
-                                       0,
-                                       d_data,
-                                       d_buffer,
-                                       d_rng_state);
-
-#define KERNEL_TEX(type, ttype, name) \
-               set_kernel_arg_mem(ckPathTraceKernel, &start_arg_index, #name);
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-
-               start_arg_index += kernel_set_args(ckPathTraceKernel,
-                                                  start_arg_index,
-                                                  d_sample,
-                                                  d_x,
-                                                  d_y,
-                                                  d_w,
-                                                  d_h,
-                                                  d_offset,
-                                                  d_stride);
-
-               enqueue_kernel(ckPathTraceKernel, d_w, d_h);
-       }
-
-       void thread_run(DeviceTask *task)
-       {
-               if(task->type == DeviceTask::FILM_CONVERT) {
-                       film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
-               }
-               else if(task->type == DeviceTask::SHADER) {
-                       shader(*task);
-               }
-               else if(task->type == DeviceTask::PATH_TRACE) {
-                       RenderTile tile;
-                       /* Keep rendering tiles until done. */
-                       while(task->acquire_tile(this, tile)) {
-                               int start_sample = tile.start_sample;
-                               int end_sample = tile.start_sample + tile.num_samples;
-
-                               for(int sample = start_sample; sample < end_sample; sample++) {
-                                       if(task->get_cancel()) {
-                                               if(task->need_finish_queue == false)
-                                                       break;
-                                       }
-
-                                       path_trace(tile, sample);
-
-                                       tile.sample = sample + 1;
-
-                                       task->update_progress(&tile);
-                               }
-
-                               /* Complete kernel execution before release tile */
-                               /* This helps in multi-device render;
-                                * The device that reaches the critical-section function
-                                * release_tile waits (stalling other devices from entering
-                                * release_tile) for all kernels to complete. If device1 (a
-                                * slow-render device) reaches release_tile first then it would
-                                * stall device2 (a fast-render device) from proceeding to render
-                                * next tile.
-                                */
-                               clFinish(cqCommandQueue);
-
-                               task->release_tile(tile);
-                       }
-               }
-       }
-};
-
-/* TODO(sergey): This is to keep tile split on OpenCL level working
- * for now, since without this view-port render does not work as it
- * should.
- *
- * Ideally it'll be done on the higher level, but we need to get ready
- * for merge rather soon, so let's keep split logic private here in
- * the file.
- */
-class SplitRenderTile : public RenderTile {
-public:
-       SplitRenderTile()
-               : RenderTile(),
-                 buffer_offset_x(0),
-                 buffer_offset_y(0),
-                 rng_state_offset_x(0),
-                 rng_state_offset_y(0),
-                 buffer_rng_state_stride(0) {}
-
-       explicit SplitRenderTile(RenderTile& tile)
-               : RenderTile(),
-                 buffer_offset_x(0),
-                 buffer_offset_y(0),
-                 rng_state_offset_x(0),
-                 rng_state_offset_y(0),
-                 buffer_rng_state_stride(0)
-       {
-               x = tile.x;
-               y = tile.y;
-               w = tile.w;
-               h = tile.h;
-               start_sample = tile.start_sample;
-               num_samples = tile.num_samples;
-               sample = tile.sample;
-               resolution = tile.resolution;
-               offset = tile.offset;
-               stride = tile.stride;
-               buffer = tile.buffer;
-               rng_state = tile.rng_state;
-               buffers = tile.buffers;
-       }
-
-       /* Split kernel is device global memory constrained;
-        * hence split kernel cant render big tile size's in
-        * one go. If the user sets a big tile size (big tile size
-        * is a term relative to the available device global memory),
-        * we split the tile further and then call path_trace on
-        * each of those split tiles. The following variables declared,
-        * assist in achieving that purpose
-        */
-       int buffer_offset_x;
-       int buffer_offset_y;
-       int rng_state_offset_x;
-       int rng_state_offset_y;
-       int buffer_rng_state_stride;
-};
-
-/* OpenCLDeviceSplitKernel's declaration/definition. */
-class OpenCLDeviceSplitKernel : public OpenCLDeviceBase
-{
-public:
-       /* Kernel declaration. */
-       cl_kernel ckPathTraceKernel_data_init;
-       cl_kernel ckPathTraceKernel_scene_intersect;
-       cl_kernel ckPathTraceKernel_lamp_emission;
-       cl_kernel ckPathTraceKernel_queue_enqueue;
-       cl_kernel ckPathTraceKernel_background_buffer_update;
-       cl_kernel ckPathTraceKernel_shader_eval;
-       cl_kernel ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao;
-       cl_kernel ckPathTraceKernel_direct_lighting;
-       cl_kernel ckPathTraceKernel_shadow_blocked;
-       cl_kernel ckPathTraceKernel_next_iteration_setup;
-       cl_kernel ckPathTraceKernel_sum_all_radiance;
-
-       /* cl_program declaration. */
-       cl_program data_init_program;
-       cl_program scene_intersect_program;
-       cl_program lamp_emission_program;
-       cl_program queue_enqueue_program;
-       cl_program background_buffer_update_program;
-       cl_program shader_eval_program;
-       cl_program holdout_emission_blurring_pathtermination_ao_program;
-       cl_program direct_lighting_program;
-       cl_program shadow_blocked_program;
-       cl_program next_iteration_setup_program;
-       cl_program sum_all_radiance_program;
-
-       /* Global memory variables [porting]; These memory is used for
-        * co-operation between different kernels; Data written by one
-        * kernel will be available to another kernel via this global
-        * memory.
-        */
-       cl_mem rng_coop;
-       cl_mem throughput_coop;
-       cl_mem L_transparent_coop;
-       cl_mem PathRadiance_coop;
-       cl_mem Ray_coop;
-       cl_mem PathState_coop;
-       cl_mem Intersection_coop;
-       cl_mem kgbuffer;  /* KernelGlobals buffer. */
-
-       /* Global buffers for ShaderData. */
-       cl_mem sd;             /* ShaderData used in the main path-iteration loop. */
-       cl_mem sd_DL_shadow;   /* ShaderData used in Direct Lighting and
-                               * shadow_blocked kernel.
-                               */
-
-       /* Global memory required for shadow blocked and accum_radiance. */
-       cl_mem BSDFEval_coop;
-       cl_mem ISLamp_coop;
-       cl_mem LightRay_coop;
-       cl_mem AOAlpha_coop;
-       cl_mem AOBSDF_coop;
-       cl_mem AOLightRay_coop;
-       cl_mem Intersection_coop_shadow;
-
-#ifdef WITH_CYCLES_DEBUG
-       /* DebugData memory */
-       cl_mem debugdata_coop;
-#endif
-
-       /* Global state array that tracks ray state. */
-       cl_mem ray_state;
-
-       /* Per sample buffers. */
-       cl_mem per_sample_output_buffers;
-
-       /* Denotes which sample each ray is being processed for. */
-       cl_mem work_array;
-
-       /* Queue */
-       cl_mem Queue_data;  /* Array of size queuesize * num_queues * sizeof(int). */
-       cl_mem Queue_index; /* Array of size num_queues * sizeof(int);
-                            * Tracks the size of each queue.
-                            */
-
-       /* Flag to make sceneintersect and lampemission kernel use queues. */
-       cl_mem use_queues_flag;
-
-       /* Amount of memory in output buffer associated with one pixel/thread. */
-       size_t per_thread_output_buffer_size;
-
-       /* Total allocatable available device memory. */
-       size_t total_allocatable_memory;
-
-       /* host version of ray_state; Used in checking host path-iteration
-        * termination.
-        */
-       char *hostRayStateArray;
-
-       /* Number of path-iterations to be done in one shot. */
-       unsigned int PathIteration_times;
-
-#ifdef __WORK_STEALING__
-       /* Work pool with respect to each work group. */
-       cl_mem work_pool_wgs;
-
-       /* Denotes the maximum work groups possible w.r.t. current tile size. */
-       unsigned int max_work_groups;
-#endif
-
-       /* clos_max value for which the kernels have been loaded currently. */
-       int current_max_closure;
-
-       /* Marked True in constructor and marked false at the end of path_trace(). */
-       bool first_tile;
-
-       OpenCLDeviceSplitKernel(DeviceInfo& info, Stats &stats, bool background_)
-       : OpenCLDeviceBase(info, stats, background_)
-       {
-               background = background_;
-
-               /* Initialize kernels. */
-               ckPathTraceKernel_data_init = NULL;
-               ckPathTraceKernel_scene_intersect = NULL;
-               ckPathTraceKernel_lamp_emission = NULL;
-               ckPathTraceKernel_background_buffer_update = NULL;
-               ckPathTraceKernel_shader_eval = NULL;
-               ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao = NULL;
-               ckPathTraceKernel_direct_lighting = NULL;
-               ckPathTraceKernel_shadow_blocked = NULL;
-               ckPathTraceKernel_next_iteration_setup = NULL;
-               ckPathTraceKernel_sum_all_radiance = NULL;
-               ckPathTraceKernel_queue_enqueue = NULL;
-
-               /* Initialize program. */
-               data_init_program = NULL;
-               scene_intersect_program = NULL;
-               lamp_emission_program = NULL;
-               queue_enqueue_program = NULL;
-               background_buffer_update_program = NULL;
-               shader_eval_program = NULL;
-               holdout_emission_blurring_pathtermination_ao_program = NULL;
-               direct_lighting_program = NULL;
-               shadow_blocked_program = NULL;
-               next_iteration_setup_program = NULL;
-               sum_all_radiance_program = NULL;
-
-               /* Initialize cl_mem variables. */
-               kgbuffer = NULL;
-               sd = NULL;
-               sd_DL_shadow = NULL;
-
-               rng_coop = NULL;
-               throughput_coop = NULL;
-               L_transparent_coop = NULL;
-               PathRadiance_coop = NULL;
-               Ray_coop = NULL;
-               PathState_coop = NULL;
-               Intersection_coop = NULL;
-               ray_state = NULL;
-
-               AOAlpha_coop = NULL;
-               AOBSDF_coop = NULL;
-               AOLightRay_coop = NULL;
-               BSDFEval_coop = NULL;
-               ISLamp_coop = NULL;
-               LightRay_coop = NULL;
-               Intersection_coop_shadow = NULL;
-
-#ifdef WITH_CYCLES_DEBUG
-               debugdata_coop = NULL;
-#endif
-
-               work_array = NULL;
-
-               /* Queue. */
-               Queue_data = NULL;
-               Queue_index = NULL;
-               use_queues_flag = NULL;
-
-               per_sample_output_buffers = NULL;
-
-               per_thread_output_buffer_size = 0;
-               hostRayStateArray = NULL;
-               PathIteration_times = PATH_ITER_INC_FACTOR;
-#ifdef __WORK_STEALING__
-               work_pool_wgs = NULL;
-               max_work_groups = 0;
-#endif
-               current_max_closure = -1;
-               first_tile = true;
-
-               /* Get device's maximum memory that can be allocated. */
-               ciErr = clGetDeviceInfo(cdDevice,
-                                       CL_DEVICE_MAX_MEM_ALLOC_SIZE,
-                                       sizeof(size_t),
-                                       &total_allocatable_memory,
-                                       NULL);
-               assert(ciErr == CL_SUCCESS);
-               if(platform_name == "AMD Accelerated Parallel Processing") {
-                       /* This value is tweak-able; AMD platform does not seem to
-                        * give maximum performance when all of CL_DEVICE_MAX_MEM_ALLOC_SIZE
-                        * is considered for further computation.
-                        */
-                       total_allocatable_memory /= 2;
-               }
-       }
-
-       /* TODO(sergey): Seems really close to load_kernel(),
-        * could it be de-duplicated?
-        */
-       bool load_split_kernel(const string& kernel_name,
-                              const string& kernel_path,
-                              const string& kernel_init_source,
-                              const string& clbin,
-                              const string& custom_kernel_build_options,
-                              cl_program *program,
-                              const string *debug_src = NULL)
-       {
-               if(!opencl_version_check()) {
-                       return false;
-               }
-
-               string cache_clbin = path_cache_get(path_join("kernels", clbin));
-
-               /* If exists already, try use it. */
-               if(path_exists(cache_clbin) && load_binary(kernel_path,
-                                                          cache_clbin,
-                                                          custom_kernel_build_options,
-                                                          program,
-                                                          debug_src))
-               {
-                       /* Kernel loaded from binary. */
-               }
-               else {
-                       /* If does not exist or loading binary failed, compile kernel. */
-                       if(!compile_kernel(kernel_name,
-                                          kernel_path,
-                                          kernel_init_source,
-                                          custom_kernel_build_options,
-                                          program,
-                                          debug_src))
-                       {
-                               return false;
-                       }
-                       /* Save binary for reuse. */
-                       if(!save_binary(program, cache_clbin)) {
-                               return false;
-                       }
-               }
-               return true;
-       }
-
-       /* Split kernel utility functions. */
-       size_t get_tex_size(const char *tex_name)
-       {
-               cl_mem ptr;
-               size_t ret_size = 0;
-               MemMap::iterator i = mem_map.find(tex_name);
-               if(i != mem_map.end()) {
-                       ptr = CL_MEM_PTR(i->second);
-                       ciErr = clGetMemObjectInfo(ptr,
-                                                  CL_MEM_SIZE,
-                                                  sizeof(ret_size),
-                                                  &ret_size,
-                                                  NULL);
-                       assert(ciErr == CL_SUCCESS);
-               }
-               return ret_size;
-       }
-
-       size_t get_shader_data_size(size_t max_closure)
-       {
-               /* ShaderData size with variable size ShaderClosure array */
-               return sizeof(ShaderData) - (sizeof(ShaderClosure) * (MAX_CLOSURE - max_closure));
-       }
-
-       /* Returns size of KernelGlobals structure associated with OpenCL. */
-       size_t get_KernelGlobals_size()
-       {
-               /* Copy dummy KernelGlobals related to OpenCL from kernel_globals.h to
-                * fetch its size.
-                */
-               typedef struct KernelGlobals {
-                       ccl_constant KernelData *data;
-#define KERNEL_TEX(type, ttype, name) \
-       ccl_global type *name;
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-                       void *sd_input;
-                       void *isect_shadow;
-               } KernelGlobals;
-
-               return sizeof(KernelGlobals);
-       }
-
-       bool load_kernels(const DeviceRequestedFeatures& requested_features)
-       {
-               /* Get Shader, bake and film_convert kernels.
-                * It'll also do verification of OpenCL actually initialized.
-                */
-               if(!OpenCLDeviceBase::load_kernels(requested_features)) {
-                       return false;
-               }
-
-               string kernel_path = path_get("kernel");
-               string kernel_md5 = path_files_md5_hash(kernel_path);
-               string device_md5;
-               string kernel_init_source;
-               string clbin;
-               string clsrc, *debug_src = NULL;
-
-               string build_options = "-D__SPLIT_KERNEL__ ";
-#ifdef __WORK_STEALING__
-               build_options += "-D__WORK_STEALING__ ";
-#endif
-               build_options += requested_features.get_build_options();
-
-               /* Set compute device build option. */
-               cl_device_type device_type;
-               ciErr = clGetDeviceInfo(cdDevice,
-                                       CL_DEVICE_TYPE,
-                                       sizeof(cl_device_type),
-                                       &device_type,
-                                       NULL);
-               assert(ciErr == CL_SUCCESS);
-               if(device_type == CL_DEVICE_TYPE_GPU) {
-                       build_options += " -D__COMPUTE_DEVICE_GPU__";
-               }
-
-#define GLUE(a, b) a ## b
-#define LOAD_KERNEL(name) \
-       do { \
-               kernel_init_source = "#include \"kernels/opencl/kernel_" #name ".cl\" // " + \
-                                    kernel_md5 + "\n"; \
-               device_md5 = device_md5_hash(build_options); \
-               clbin = string_printf("cycles_kernel_%s_%s_" #name ".clbin", \
-                                     device_md5.c_str(), kernel_md5.c_str()); \
-               if(opencl_kernel_use_debug()) { \
-                       clsrc = string_printf("cycles_kernel_%s_%s_" #name ".cl", \
-                                             device_md5.c_str(), kernel_md5.c_str()); \
-                       clsrc = path_cache_get(path_join("kernels", clsrc)); \
-                       debug_src = &clsrc; \
-               } \
-               if(!load_split_kernel(#name, \
-                                     kernel_path, \
-                                     kernel_init_source, \
-                                     clbin, \
-                                     build_options, \
-                                     &GLUE(name, _program), \
-                                     debug_src)) \
-               { \
-                       fprintf(stderr, "Faled to compile %s\n", #name); \
-                       return false; \
-               } \
-       } while(false)
-
-               LOAD_KERNEL(data_init);
-               LOAD_KERNEL(scene_intersect);
-               LOAD_KERNEL(lamp_emission);
-               LOAD_KERNEL(queue_enqueue);
-               LOAD_KERNEL(background_buffer_update);
-               LOAD_KERNEL(shader_eval);
-               LOAD_KERNEL(holdout_emission_blurring_pathtermination_ao);
-               LOAD_KERNEL(direct_lighting);
-               LOAD_KERNEL(shadow_blocked);
-               LOAD_KERNEL(next_iteration_setup);
-               LOAD_KERNEL(sum_all_radiance);
-
-#undef LOAD_KERNEL
-
-#define FIND_KERNEL(name) \
-       do { \
-               GLUE(ckPathTraceKernel_, name) = \
-                       clCreateKernel(GLUE(name, _program), \
-                                      "kernel_ocl_path_trace_"  #name, &ciErr); \
-               if(opencl_error(ciErr)) { \
-                       fprintf(stderr,"Missing kernel kernel_ocl_path_trace_%s\n", #name); \
-                       return false; \
-               } \
-       } while(false)
-
-               FIND_KERNEL(data_init);
-               FIND_KERNEL(scene_intersect);
-               FIND_KERNEL(lamp_emission);
-               FIND_KERNEL(queue_enqueue);
-               FIND_KERNEL(background_buffer_update);
-               FIND_KERNEL(shader_eval);
-               FIND_KERNEL(holdout_emission_blurring_pathtermination_ao);
-               FIND_KERNEL(direct_lighting);
-               FIND_KERNEL(shadow_blocked);
-               FIND_KERNEL(next_iteration_setup);
-               FIND_KERNEL(sum_all_radiance);
-#undef FIND_KERNEL
-#undef GLUE
-
-               current_max_closure = requested_features.max_closure;
-
-               return true;
-       }
-
-       ~OpenCLDeviceSplitKernel()
-       {
-               task_pool.stop();
-
-               /* Release kernels */
-               release_kernel_safe(ckPathTraceKernel_data_init);
-               release_kernel_safe(ckPathTraceKernel_scene_intersect);
-               release_kernel_safe(ckPathTraceKernel_lamp_emission);
-               release_kernel_safe(ckPathTraceKernel_queue_enqueue);
-               release_kernel_safe(ckPathTraceKernel_background_buffer_update);
-               release_kernel_safe(ckPathTraceKernel_shader_eval);
-               release_kernel_safe(ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao);
-               release_kernel_safe(ckPathTraceKernel_direct_lighting);
-               release_kernel_safe(ckPathTraceKernel_shadow_blocked);
-               release_kernel_safe(ckPathTraceKernel_next_iteration_setup);
-               release_kernel_safe(ckPathTraceKernel_sum_all_radiance);
-
-               /* Release global memory */
-               release_mem_object_safe(rng_coop);
-               release_mem_object_safe(throughput_coop);
-               release_mem_object_safe(L_transparent_coop);
-               release_mem_object_safe(PathRadiance_coop);
-               release_mem_object_safe(Ray_coop);
-               release_mem_object_safe(PathState_coop);
-               release_mem_object_safe(Intersection_coop);
-               release_mem_object_safe(kgbuffer);
-               release_mem_object_safe(sd);
-               release_mem_object_safe(sd_DL_shadow);
-               release_mem_object_safe(ray_state);
-               release_mem_object_safe(AOAlpha_coop);
-               release_mem_object_safe(AOBSDF_coop);
-               release_mem_object_safe(AOLightRay_coop);
-               release_mem_object_safe(BSDFEval_coop);
-               release_mem_object_safe(ISLamp_coop);
-               release_mem_object_safe(LightRay_coop);
-               release_mem_object_safe(Intersection_coop_shadow);
-#ifdef WITH_CYCLES_DEBUG
-               release_mem_object_safe(debugdata_coop);
-#endif
-               release_mem_object_safe(use_queues_flag);
-               release_mem_object_safe(Queue_data);
-               release_mem_object_safe(Queue_index);
-               release_mem_object_safe(work_array);
-#ifdef __WORK_STEALING__
-               release_mem_object_safe(work_pool_wgs);
-#endif
-               release_mem_object_safe(per_sample_output_buffers);
-
-               /* Release programs */
-               release_program_safe(data_init_program);
-               release_program_safe(scene_intersect_program);
-               release_program_safe(lamp_emission_program);
-               release_program_safe(queue_enqueue_program);
-               release_program_safe(background_buffer_update_program);
-               release_program_safe(shader_eval_program);
-               release_program_safe(holdout_emission_blurring_pathtermination_ao_program);
-               release_program_safe(direct_lighting_program);
-               release_program_safe(shadow_blocked_program);
-               release_program_safe(next_iteration_setup_program);
-               release_program_safe(sum_all_radiance_program);
-
-               if(hostRayStateArray != NULL) {
-                       free(hostRayStateArray);
-               }
-       }
-
-       void path_trace(DeviceTask *task,
-                       SplitRenderTile& rtile,
-                       int2 max_render_feasible_tile_size)
-       {
-               /* cast arguments to cl types */
-               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
-               cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
-               cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
-               cl_int d_x = rtile.x;
-               cl_int d_y = rtile.y;
-               cl_int d_w = rtile.w;
-               cl_int d_h = rtile.h;
-               cl_int d_offset = rtile.offset;
-               cl_int d_stride = rtile.stride;
-
-               /* Make sure that set render feasible tile size is a multiple of local
-                * work size dimensions.
-                */
-               assert(max_render_feasible_tile_size.x % SPLIT_KERNEL_LOCAL_SIZE_X == 0);
-               assert(max_render_feasible_tile_size.y % SPLIT_KERNEL_LOCAL_SIZE_Y == 0);
-
-               size_t global_size[2];
-               size_t local_size[2] = {SPLIT_KERNEL_LOCAL_SIZE_X,
-                                       SPLIT_KERNEL_LOCAL_SIZE_Y};
-
-               /* Set the range of samples to be processed for every ray in
-                * path-regeneration logic.
-                */
-               cl_int start_sample = rtile.start_sample;
-               cl_int end_sample = rtile.start_sample + rtile.num_samples;
-               cl_int num_samples = rtile.num_samples;
-
-#ifdef __WORK_STEALING__
-               global_size[0] = (((d_w - 1) / local_size[0]) + 1) * local_size[0];
-               global_size[1] = (((d_h - 1) / local_size[1]) + 1) * local_size[1];
-               unsigned int num_parallel_samples = 1;
-#else
-               global_size[1] = (((d_h - 1) / local_size[1]) + 1) * local_size[1];
-               unsigned int num_threads = max_render_feasible_tile_size.x *
-                                          max_render_feasible_tile_size.y;
-               unsigned int num_tile_columns_possible = num_threads / global_size[1];
-               /* Estimate number of parallel samples that can be
-                * processed in parallel.
-                */
-               unsigned int num_parallel_samples = min(num_tile_columns_possible / d_w,
-                                                       rtile.num_samples);
-               /* Wavefront size in AMD is 64.
-                * TODO(sergey): What about other platforms?
-                */
-               if(num_parallel_samples >= 64) {
-                       /* TODO(sergey): Could use generic round-up here. */
-                       num_parallel_samples = (num_parallel_samples / 64) * 64;
-               }
-               assert(num_parallel_samples != 0);
-
-               global_size[0] = d_w * num_parallel_samples;
-#endif  /* __WORK_STEALING__ */
-
-               assert(global_size[0] * global_size[1] <=
-                      max_render_feasible_tile_size.x * max_render_feasible_tile_size.y);
-
-               /* Allocate all required global memory once. */
-               if(first_tile) {
-                       size_t num_global_elements = max_render_feasible_tile_size.x *
-                                                    max_render_feasible_tile_size.y;
-                       /* TODO(sergey): This will actually over-allocate if
-                        * particular kernel does not support multiclosure.
-                        */
-                       size_t shaderdata_size = get_shader_data_size(current_max_closure);
-
-#ifdef __WORK_STEALING__
-                       /* Calculate max groups */
-                       size_t max_global_size[2];
-                       size_t tile_x = max_render_feasible_tile_size.x;
-                       size_t tile_y = max_render_feasible_tile_size.y;
-                       max_global_size[0] = (((tile_x - 1) / local_size[0]) + 1) * local_size[0];
-                       max_global_size[1] = (((tile_y - 1) / local_size[1]) + 1) * local_size[1];
-                       max_work_groups = (max_global_size[0] * max_global_size[1]) /
-                                         (local_size[0] * local_size[1]);
-                       /* Allocate work_pool_wgs memory. */
-                       work_pool_wgs = mem_alloc(max_work_groups * sizeof(unsigned int));
-#endif  /* __WORK_STEALING__ */
-
-                       /* Allocate queue_index memory only once. */
-                       Queue_index = mem_alloc(NUM_QUEUES * sizeof(int));
-                       use_queues_flag = mem_alloc(sizeof(char));
-                       kgbuffer = mem_alloc(get_KernelGlobals_size());
-
-                       /* Create global buffers for ShaderData. */
-                       sd = mem_alloc(num_global_elements * shaderdata_size);
-                       sd_DL_shadow = mem_alloc(num_global_elements * 2 * shaderdata_size);
-
-                       /* Creation of global memory buffers which are shared among
-                        * the kernels.
-                        */
-                       rng_coop = mem_alloc(num_global_elements * sizeof(RNG));
-                       throughput_coop = mem_alloc(num_global_elements * sizeof(float3));
-                       L_transparent_coop = mem_alloc(num_global_elements * sizeof(float));
-                       PathRadiance_coop = mem_alloc(num_global_elements * sizeof(PathRadiance));
-                       Ray_coop = mem_alloc(num_global_elements * sizeof(Ray));
-                       PathState_coop = mem_alloc(num_global_elements * sizeof(PathState));
-                       Intersection_coop = mem_alloc(num_global_elements * sizeof(Intersection));
-                       AOAlpha_coop = mem_alloc(num_global_elements * sizeof(float3));
-                       AOBSDF_coop = mem_alloc(num_global_elements * sizeof(float3));
-                       AOLightRay_coop = mem_alloc(num_global_elements * sizeof(Ray));
-                       BSDFEval_coop = mem_alloc(num_global_elements * sizeof(BsdfEval));
-                       ISLamp_coop = mem_alloc(num_global_elements * sizeof(int));
-                       LightRay_coop = mem_alloc(num_global_elements * sizeof(Ray));
-                       Intersection_coop_shadow = mem_alloc(2 * num_global_elements * sizeof(Intersection));
-
-#ifdef WITH_CYCLES_DEBUG
-                       debugdata_coop = mem_alloc(num_global_elements * sizeof(DebugData));
-#endif
-
-                       ray_state = mem_alloc(num_global_elements * sizeof(char));
-
-                       hostRayStateArray = (char *)calloc(num_global_elements, sizeof(char));
-                       assert(hostRayStateArray != NULL && "Can't create hostRayStateArray memory");
-
-                       Queue_data = mem_alloc(num_global_elements * (NUM_QUEUES * sizeof(int)+sizeof(int)));
-                       work_array = mem_alloc(num_global_elements * sizeof(unsigned int));
-                       per_sample_output_buffers = mem_alloc(num_global_elements *
-                                                             per_thread_output_buffer_size);
-               }
-
-               cl_int dQueue_size = global_size[0] * global_size[1];
-
-               cl_uint start_arg_index =
-                       kernel_set_args(ckPathTraceKernel_data_init,
-                                       0,
-                                       kgbuffer,
-                                       sd_DL_shadow,
-                                       d_data,
-                                       per_sample_output_buffers,
-                                       d_rng_state,
-                                       rng_coop,
-                                       throughput_coop,
-                                       L_transparent_coop,
-                                       PathRadiance_coop,
-                                       Ray_coop,
-                                       PathState_coop,
-                                       Intersection_coop_shadow,
-                                       ray_state);
-
-/* TODO(sergey): Avoid map lookup here. */
-#define KERNEL_TEX(type, ttype, name) \
-       set_kernel_arg_mem(ckPathTraceKernel_data_init, &start_arg_index, #name);
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-
-               start_arg_index +=
-                       kernel_set_args(ckPathTraceKernel_data_init,
-                                       start_arg_index,
-                                       start_sample,
-                                       d_x,
-                                       d_y,
-                                       d_w,
-                                       d_h,
-                                       d_offset,
-                                       d_stride,
-                                       rtile.rng_state_offset_x,
-                                       rtile.rng_state_offset_y,
-                                       rtile.buffer_rng_state_stride,
-                                       Queue_data,
-                                       Queue_index,
-                                       dQueue_size,
-                                       use_queues_flag,
-                                       work_array,
-#ifdef __WORK_STEALING__
-                                       work_pool_wgs,
-                                       num_samples,
-#endif
-#ifdef WITH_CYCLES_DEBUG
-                                       debugdata_coop,
-#endif
-                                       num_parallel_samples);
-
-               kernel_set_args(ckPathTraceKernel_scene_intersect,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               rng_coop,
-                               Ray_coop,
-                               PathState_coop,
-                               Intersection_coop,
-                               ray_state,
-                               d_w,
-                               d_h,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size,
-                               use_queues_flag,
-#ifdef WITH_CYCLES_DEBUG
-                               debugdata_coop,
-#endif
-                               num_parallel_samples);
-
-               kernel_set_args(ckPathTraceKernel_lamp_emission,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               throughput_coop,
-                               PathRadiance_coop,
-                               Ray_coop,
-                               PathState_coop,
-                               Intersection_coop,
-                               ray_state,
-                               d_w,
-                               d_h,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size,
-                               use_queues_flag,
-                               num_parallel_samples);
-
-               kernel_set_args(ckPathTraceKernel_queue_enqueue,
-                               0,
-                               Queue_data,
-                               Queue_index,
-                               ray_state,
-                               dQueue_size);
-
-               kernel_set_args(ckPathTraceKernel_background_buffer_update,
-                                0,
-                                kgbuffer,
-                                d_data,
-                                per_sample_output_buffers,
-                                d_rng_state,
-                                rng_coop,
-                                throughput_coop,
-                                PathRadiance_coop,
-                                Ray_coop,
-                                PathState_coop,
-                                L_transparent_coop,
-                                ray_state,
-                                d_w,
-                                d_h,
-                                d_x,
-                                d_y,
-                                d_stride,
-                                rtile.rng_state_offset_x,
-                                rtile.rng_state_offset_y,
-                                rtile.buffer_rng_state_stride,
-                                work_array,
-                                Queue_data,
-                                Queue_index,
-                                dQueue_size,
-                                end_sample,
-                                start_sample,
-#ifdef __WORK_STEALING__
-                                work_pool_wgs,
-                                num_samples,
-#endif
-#ifdef WITH_CYCLES_DEBUG
-                                debugdata_coop,
-#endif
-                                num_parallel_samples);
-
-               kernel_set_args(ckPathTraceKernel_shader_eval,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               sd,
-                               rng_coop,
-                               Ray_coop,
-                               PathState_coop,
-                               Intersection_coop,
-                               ray_state,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size);
-
-               kernel_set_args(ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               sd,
-                               per_sample_output_buffers,
-                               rng_coop,
-                               throughput_coop,
-                               L_transparent_coop,
-                               PathRadiance_coop,
-                               PathState_coop,
-                               Intersection_coop,
-                               AOAlpha_coop,
-                               AOBSDF_coop,
-                               AOLightRay_coop,
-                               d_w,
-                               d_h,
-                               d_x,
-                               d_y,
-                               d_stride,
-                               ray_state,
-                               work_array,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size,
-#ifdef __WORK_STEALING__
-                               start_sample,
-#endif
-                               num_parallel_samples);
-
-               kernel_set_args(ckPathTraceKernel_direct_lighting,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               sd,
-                               rng_coop,
-                               PathState_coop,
-                               ISLamp_coop,
-                               LightRay_coop,
-                               BSDFEval_coop,
-                               ray_state,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size);
-
-               kernel_set_args(ckPathTraceKernel_shadow_blocked,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               PathState_coop,
-                               LightRay_coop,
-                               AOLightRay_coop,
-                               ray_state,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size);
-
-               kernel_set_args(ckPathTraceKernel_next_iteration_setup,
-                               0,
-                               kgbuffer,
-                               d_data,
-                               sd,
-                               rng_coop,
-                               throughput_coop,
-                               PathRadiance_coop,
-                               Ray_coop,
-                               PathState_coop,
-                               LightRay_coop,
-                               ISLamp_coop,
-                               BSDFEval_coop,
-                               AOLightRay_coop,
-                               AOBSDF_coop,
-                               AOAlpha_coop,
-                               ray_state,
-                               Queue_data,
-                               Queue_index,
-                               dQueue_size,
-                               use_queues_flag);
-
-               kernel_set_args(ckPathTraceKernel_sum_all_radiance,
-                               0,
-                               d_data,
-                               d_buffer,
-                               per_sample_output_buffers,
-                               num_parallel_samples,
-                               d_w,
-                               d_h,
-                               d_stride,
-                               rtile.buffer_offset_x,
-                               rtile.buffer_offset_y,
-                               rtile.buffer_rng_state_stride,
-                               start_sample);
-
-               /* Macro for Enqueuing split kernels. */
-#define GLUE(a, b) a ## b
-#define ENQUEUE_SPLIT_KERNEL(kernelName, globalSize, localSize) \
-               { \
-                       ciErr = clEnqueueNDRangeKernel(cqCommandQueue, \
-                                                      GLUE(ckPathTraceKernel_, \
-                                                           kernelName), \
-                                                      2, \
-                                                      NULL, \
-                                                      globalSize, \
-                                                      localSize, \
-                                                      0, \
-                                                      NULL, \
-                                                      NULL); \
-                       opencl_assert_err(ciErr, "clEnqueueNDRangeKernel"); \
-                       if(ciErr != CL_SUCCESS) { \
-                               string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()", \
-                                                              clewErrorString(ciErr)); \
-                               opencl_error(message); \
-                               return; \
-                       } \
-               } (void) 0
-
-               /* Enqueue ckPathTraceKernel_data_init kernel. */
-               ENQUEUE_SPLIT_KERNEL(data_init, global_size, local_size);
-               bool activeRaysAvailable = true;
-
-               /* Record number of time host intervention has been made */
-               unsigned int numHostIntervention = 0;
-               unsigned int numNextPathIterTimes = PathIteration_times;
-               bool canceled = false;
-               while(activeRaysAvailable) {
-                       /* Twice the global work size of other kernels for
-                        * ckPathTraceKernel_shadow_blocked_direct_lighting. */
-                       size_t global_size_shadow_blocked[2];
-                       global_size_shadow_blocked[0] = global_size[0] * 2;
-                       global_size_shadow_blocked[1] = global_size[1];
-
-                       /* Do path-iteration in host [Enqueue Path-iteration kernels. */
-                       for(int PathIter = 0; PathIter < PathIteration_times; PathIter++) {
-                               ENQUEUE_SPLIT_KERNEL(scene_intersect, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(lamp_emission, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(background_buffer_update, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(shader_eval, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(holdout_emission_blurring_pathtermination_ao, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
-                               ENQUEUE_SPLIT_KERNEL(shadow_blocked, global_size_shadow_blocked, local_size);
-                               ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size);
-                               if(task->get_cancel()) {
-                                       canceled = true;
-                                       break;
-                               }
-                       }
-
-                       /* Read ray-state into Host memory to decide if we should exit
-                        * path-iteration in host.
-                        */
-                       ciErr = clEnqueueReadBuffer(cqCommandQueue,
-                                                   ray_state,
-                                                   CL_TRUE,
-                                                   0,
-                                                   global_size[0] * global_size[1] * sizeof(char),
-                                                   hostRayStateArray,
-                                                   0,
-                                                   NULL,
-                                                   NULL);
-                       assert(ciErr == CL_SUCCESS);
-
-                       activeRaysAvailable = false;
-
-                       for(int rayStateIter = 0;
-                           rayStateIter < global_size[0] * global_size[1];
-                           ++rayStateIter)
-                       {
-                               if(int8_t(hostRayStateArray[rayStateIter]) != RAY_INACTIVE) {
-                                       /* Not all rays are RAY_INACTIVE. */
-                                       activeRaysAvailable = true;
-                                       break;
-                               }
-                       }
-
-                       if(activeRaysAvailable) {
-                               numHostIntervention++;
-                               PathIteration_times = PATH_ITER_INC_FACTOR;
-                               /* Host intervention done before all rays become RAY_INACTIVE;
-                                * Set do more initial iterations for the next tile.
-                                */
-                               numNextPathIterTimes += PATH_ITER_INC_FACTOR;
-                       }
-                       if(task->get_cancel()) {
-                               canceled = true;
-                               break;
-                       }
-               }
-
-               /* Execute SumALLRadiance kernel to accumulate radiance calculated in
-                * per_sample_output_buffers into RenderTile's output buffer.
-                */
-               if (!canceled) {
-                       size_t sum_all_radiance_local_size[2] = {16, 16};
-                       size_t sum_all_radiance_global_size[2];
-                       sum_all_radiance_global_size[0] =
-                               (((d_w - 1) / sum_all_radiance_local_size[0]) + 1) *
-                               sum_all_radiance_local_size[0];
-                       sum_all_radiance_global_size[1] =
-                               (((d_h - 1) / sum_all_radiance_local_size[1]) + 1) *
-                               sum_all_radiance_local_size[1];
-                       ENQUEUE_SPLIT_KERNEL(sum_all_radiance,
-                                            sum_all_radiance_global_size,
-                                            sum_all_radiance_local_size);
-               }
-
-#undef ENQUEUE_SPLIT_KERNEL
-#undef GLUE
-
-               if(numHostIntervention == 0) {
-                       /* This means that we are executing kernel more than required
-                        * Must avoid this for the next sample/tile.
-                        */
-                       PathIteration_times = ((numNextPathIterTimes - PATH_ITER_INC_FACTOR) <= 0) ?
-                       PATH_ITER_INC_FACTOR : numNextPathIterTimes - PATH_ITER_INC_FACTOR;
-               }
-               else {
-                       /* Number of path-iterations done for this tile is set as
-                        * Initial path-iteration times for the next tile
-                        */
-                       PathIteration_times = numNextPathIterTimes;
-               }
-
-               first_tile = false;
-       }
-
-       /* Calculates the amount of memory that has to be always
-        * allocated in order for the split kernel to function.
-        * This memory is tile/scene-property invariant (meaning,
-        * the value returned by this function does not depend
-        * on the user set tile size or scene properties.
-        */
-       size_t get_invariable_mem_allocated()
-       {
-               size_t total_invariable_mem_allocated = 0;
-               size_t KernelGlobals_size = 0;
-
-               KernelGlobals_size = get_KernelGlobals_size();
-
-               total_invariable_mem_allocated += KernelGlobals_size; /* KernelGlobals size */
-               total_invariable_mem_allocated += NUM_QUEUES * sizeof(unsigned int); /* Queue index size */
-               total_invariable_mem_allocated += sizeof(char); /* use_queues_flag size */
-
-               return total_invariable_mem_allocated;
-       }
-
-       /* Calculate the memory that has-to-be/has-been allocated for
-        * the split kernel to function.
-        */
-       size_t get_tile_specific_mem_allocated(const int2 tile_size)
-       {
-               size_t tile_specific_mem_allocated = 0;
-
-               /* Get required tile info */
-               unsigned int user_set_tile_w = tile_size.x;
-               unsigned int user_set_tile_h = tile_size.y;
-
-#ifdef __WORK_STEALING__
-               /* Calculate memory to be allocated for work_pools in
-                * case of work_stealing.
-                */
-               size_t max_global_size[2];
-               size_t max_num_work_pools = 0;
-               max_global_size[0] =
-                       (((user_set_tile_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_X;
-               max_global_size[1] =
-                       (((user_set_tile_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-               max_num_work_pools =
-                       (max_global_size[0] * max_global_size[1]) /
-                       (SPLIT_KERNEL_LOCAL_SIZE_X * SPLIT_KERNEL_LOCAL_SIZE_Y);
-               tile_specific_mem_allocated += max_num_work_pools * sizeof(unsigned int);
-#endif
-
-               tile_specific_mem_allocated +=
-                       user_set_tile_w * user_set_tile_h * per_thread_output_buffer_size;
-               tile_specific_mem_allocated +=
-                       user_set_tile_w * user_set_tile_h * sizeof(RNG);
-
-               return tile_specific_mem_allocated;
-       }
-
-       /* Calculates the texture memories and KernelData (d_data) memory
-        * that has been allocated.
-        */
-       size_t get_scene_specific_mem_allocated(cl_mem d_data)
-       {
-               size_t scene_specific_mem_allocated = 0;
-               /* Calculate texture memories. */
-#define KERNEL_TEX(type, ttype, name) \
-       scene_specific_mem_allocated += get_tex_size(#name);
-#include "kernel_textures.h"
-#undef KERNEL_TEX
-               size_t d_data_size;
-               ciErr = clGetMemObjectInfo(d_data,
-                                          CL_MEM_SIZE,
-                                          sizeof(d_data_size),
-                                          &d_data_size,
-                                          NULL);
-               assert(ciErr == CL_SUCCESS && "Can't get d_data mem object info");
-               scene_specific_mem_allocated += d_data_size;
-               return scene_specific_mem_allocated;
-       }
-
-       /* Calculate the memory required for one thread in split kernel. */
-       size_t get_per_thread_memory()
-       {
-               size_t shaderdata_size = 0;
-               /* TODO(sergey): This will actually over-allocate if
-                * particular kernel does not support multiclosure.
-                */
-               shaderdata_size = get_shader_data_size(current_max_closure);
-               size_t retval = sizeof(RNG)
-                       + sizeof(float3)          /* Throughput size */
-                       + sizeof(float)           /* L transparent size */
-                       + sizeof(char)            /* Ray state size */
-                       + sizeof(unsigned int)    /* Work element size */
-                       + sizeof(int)             /* ISLamp_size */
-                       + sizeof(PathRadiance) + sizeof(Ray) + sizeof(PathState)
-                       + sizeof(Intersection)    /* Overall isect */
-                       + sizeof(Intersection)    /* Instersection_coop_AO */
-                       + sizeof(Intersection)    /* Intersection coop DL */
-                       + shaderdata_size         /* Overall ShaderData */
-                       + (shaderdata_size * 2)   /* ShaderData : DL and shadow */
-                       + sizeof(Ray) + sizeof(BsdfEval)
-                       + sizeof(float3)          /* AOAlpha size */
-                       + sizeof(float3)          /* AOBSDF size */
-                       + sizeof(Ray)
-                       + (sizeof(int) * NUM_QUEUES)
-                       + per_thread_output_buffer_size;
-               return retval;
-       }
-
-       /* Considers the total memory available in the device and
-        * and returns the maximum global work size possible.
-        */
-       size_t get_feasible_global_work_size(int2 tile_size, cl_mem d_data)
-       {
-               /* Calculate invariably allocated memory. */
-               size_t invariable_mem_allocated = get_invariable_mem_allocated();
-               /* Calculate tile specific allocated memory. */
-               size_t tile_specific_mem_allocated =
-                       get_tile_specific_mem_allocated(tile_size);
-               /* Calculate scene specific allocated memory. */
-               size_t scene_specific_mem_allocated =
-                       get_scene_specific_mem_allocated(d_data);
-               /* Calculate total memory available for the threads in global work size. */
-               size_t available_memory = total_allocatable_memory
-                       - invariable_mem_allocated
-                       - tile_specific_mem_allocated
-                       - scene_specific_mem_allocated
-                       - DATA_ALLOCATION_MEM_FACTOR;
-               size_t per_thread_memory_required = get_per_thread_memory();
-               return (available_memory / per_thread_memory_required);
-       }
-
-       /* Checks if the device has enough memory to render the whole tile;
-        * If not, we should split single tile into multiple tiles of small size
-        * and process them all.
-        */
-       bool need_to_split_tile(unsigned int d_w,
-                               unsigned int d_h,
-                               int2 max_render_feasible_tile_size)
-       {
-               size_t global_size_estimate[2];
-               /* TODO(sergey): Such round-ups are in quite few places, need to replace
-                * them with an utility macro.
-                */
-               global_size_estimate[0] =
-                       (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_X;
-               global_size_estimate[1] =
-                       (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-               if((global_size_estimate[0] * global_size_estimate[1]) >
-                  (max_render_feasible_tile_size.x * max_render_feasible_tile_size.y))
-               {
-                       return true;
-               }
-               else {
-                       return false;
-               }
-       }
-
-       /* Considers the scene properties, global memory available in the device
-        * and returns a rectanglular tile dimension (approx the maximum)
-        * that should render on split kernel.
-        */
-       int2 get_max_render_feasible_tile_size(size_t feasible_global_work_size)
-       {
-               int2 max_render_feasible_tile_size;
-               int square_root_val = (int)sqrt(feasible_global_work_size);
-               max_render_feasible_tile_size.x = square_root_val;
-               max_render_feasible_tile_size.y = square_root_val;
-               /* Ciel round-off max_render_feasible_tile_size. */
-               int2 ceil_render_feasible_tile_size;
-               ceil_render_feasible_tile_size.x =
-                       (((max_render_feasible_tile_size.x - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_X;
-               ceil_render_feasible_tile_size.y =
-                       (((max_render_feasible_tile_size.y - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-               if(ceil_render_feasible_tile_size.x * ceil_render_feasible_tile_size.y <=
-                  feasible_global_work_size)
-               {
-                       return ceil_render_feasible_tile_size;
-               }
-               /* Floor round-off max_render_feasible_tile_size. */
-               int2 floor_render_feasible_tile_size;
-               floor_render_feasible_tile_size.x =
-                       (max_render_feasible_tile_size.x / SPLIT_KERNEL_LOCAL_SIZE_X) *
-                       SPLIT_KERNEL_LOCAL_SIZE_X;
-               floor_render_feasible_tile_size.y =
-                       (max_render_feasible_tile_size.y / SPLIT_KERNEL_LOCAL_SIZE_Y) *
-                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-               return floor_render_feasible_tile_size;
-       }
-
-       /* Try splitting the current tile into multiple smaller
-        * almost-square-tiles.
-        */
-       int2 get_split_tile_size(RenderTile rtile,
-                                int2 max_render_feasible_tile_size)
-       {
-               int2 split_tile_size;
-               int num_global_threads = max_render_feasible_tile_size.x *
-                                        max_render_feasible_tile_size.y;
-               int d_w = rtile.w;
-               int d_h = rtile.h;
-               /* Ceil round off d_w and d_h */
-               d_w = (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_X;
-               d_h = (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-               while(d_w * d_h > num_global_threads) {
-                       /* Halve the longer dimension. */
-                       if(d_w >= d_h) {
-                               d_w = d_w / 2;
-                               d_w = (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                                       SPLIT_KERNEL_LOCAL_SIZE_X;
-                       }
-                       else {
-                               d_h = d_h / 2;
-                               d_h = (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                                       SPLIT_KERNEL_LOCAL_SIZE_Y;
-                       }
-               }
-               split_tile_size.x = d_w;
-               split_tile_size.y = d_h;
-               return split_tile_size;
-       }
-
-       /* Splits existing tile into multiple tiles of tile size split_tile_size. */
-       vector<SplitRenderTile> split_tiles(RenderTile rtile, int2 split_tile_size)
-       {
-               vector<SplitRenderTile> to_path_trace_rtile;
-               int d_w = rtile.w;
-               int d_h = rtile.h;
-               int num_tiles_x = (((d_w - 1) / split_tile_size.x) + 1);
-               int num_tiles_y = (((d_h - 1) / split_tile_size.y) + 1);
-               /* Buffer and rng_state offset calc. */
-               size_t offset_index = rtile.offset + (rtile.x + rtile.y * rtile.stride);
-               size_t offset_x = offset_index % rtile.stride;
-               size_t offset_y = offset_index / rtile.stride;
-               /* Resize to_path_trace_rtile. */
-               to_path_trace_rtile.resize(num_tiles_x * num_tiles_y);
-               for(int tile_iter_y = 0; tile_iter_y < num_tiles_y; tile_iter_y++) {
-                       for(int tile_iter_x = 0; tile_iter_x < num_tiles_x; tile_iter_x++) {
-                               int rtile_index = tile_iter_y * num_tiles_x + tile_iter_x;
-                               to_path_trace_rtile[rtile_index].rng_state_offset_x = offset_x + tile_iter_x * split_tile_size.x;
-                               to_path_trace_rtile[rtile_index].rng_state_offset_y = offset_y + tile_iter_y * split_tile_size.y;
-                               to_path_trace_rtile[rtile_index].buffer_offset_x = offset_x + tile_iter_x * split_tile_size.x;
-                               to_path_trace_rtile[rtile_index].buffer_offset_y = offset_y + tile_iter_y * split_tile_size.y;
-                               to_path_trace_rtile[rtile_index].start_sample = rtile.start_sample;
-                               to_path_trace_rtile[rtile_index].num_samples = rtile.num_samples;
-                               to_path_trace_rtile[rtile_index].sample = rtile.sample;
-                               to_path_trace_rtile[rtile_index].resolution = rtile.resolution;
-                               to_path_trace_rtile[rtile_index].offset = rtile.offset;
-                               to_path_trace_rtile[rtile_index].buffers = rtile.buffers;
-                               to_path_trace_rtile[rtile_index].buffer = rtile.buffer;
-                               to_path_trace_rtile[rtile_index].rng_state = rtile.rng_state;
-                               to_path_trace_rtile[rtile_index].x = rtile.x + (tile_iter_x * split_tile_size.x);
-                               to_path_trace_rtile[rtile_index].y = rtile.y + (tile_iter_y * split_tile_size.y);
-                               to_path_trace_rtile[rtile_index].buffer_rng_state_stride = rtile.stride;
-                               /* Fill width and height of the new render tile. */
-                               to_path_trace_rtile[rtile_index].w = (tile_iter_x == (num_tiles_x - 1)) ?
-                                       (d_w - (tile_iter_x * split_tile_size.x)) /* Border tile */
-                                       : split_tile_size.x;
-                               to_path_trace_rtile[rtile_index].h = (tile_iter_y == (num_tiles_y - 1)) ?
-                                       (d_h - (tile_iter_y * split_tile_size.y)) /* Border tile */
-                                       : split_tile_size.y;
-                               to_path_trace_rtile[rtile_index].stride = to_path_trace_rtile[rtile_index].w;
-                       }
-               }
-               return to_path_trace_rtile;
-       }
-
-       void thread_run(DeviceTask *task)
-       {
-               if(task->type == DeviceTask::FILM_CONVERT) {
-                       film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
-               }
-               else if(task->type == DeviceTask::SHADER) {
-                       shader(*task);
-               }
-               else if(task->type == DeviceTask::PATH_TRACE) {
-                       RenderTile tile;
-                       bool initialize_data_and_check_render_feasibility = false;
-                       bool need_to_split_tiles_further = false;
-                       int2 max_render_feasible_tile_size;
-                       size_t feasible_global_work_size;
-                       const int2 tile_size = task->requested_tile_size;
-                       /* Keep rendering tiles until done. */
-                       while(task->acquire_tile(this, tile)) {
-                               if(!initialize_data_and_check_render_feasibility) {
-                                       /* Initialize data. */
-                                       /* Calculate per_thread_output_buffer_size. */
-                                       size_t output_buffer_size = 0;
-                                       ciErr = clGetMemObjectInfo((cl_mem)tile.buffer,
-                                                                  CL_MEM_SIZE,
-                                                                  sizeof(output_buffer_size),
-                                                                  &output_buffer_size,
-                                                                  NULL);
-                                       assert(ciErr == CL_SUCCESS && "Can't get tile.buffer mem object info");
-                                       /* This value is different when running on AMD and NV. */
-                                       if(background) {
-                                               /* In offline render the number of buffer elements
-                                                * associated with tile.buffer is the current tile size.
-                                                */
-                                               per_thread_output_buffer_size =
-                                                       output_buffer_size / (tile.w * tile.h);
-                                       }
-                                       else {
-                                               /* interactive rendering, unlike offline render, the number of buffer elements
-                                                * associated with tile.buffer is the entire viewport size.
-                                                */
-                                               per_thread_output_buffer_size =
-                                                       output_buffer_size / (tile.buffers->params.width *
-                                                                             tile.buffers->params.height);
-                                       }
-                                       /* Check render feasibility. */
-                                       feasible_global_work_size = get_feasible_global_work_size(
-                                               tile_size,
-                                               CL_MEM_PTR(const_mem_map["__data"]->device_pointer));
-                                       max_render_feasible_tile_size =
-                                               get_max_render_feasible_tile_size(
-                                                       feasible_global_work_size);
-                                       need_to_split_tiles_further =
-                                               need_to_split_tile(tile_size.x,
-                                                                  tile_size.y,
-                                                                  max_render_feasible_tile_size);
-                                       initialize_data_and_check_render_feasibility = true;
-                               }
-                               if(need_to_split_tiles_further) {
-                                       int2 split_tile_size =
-                                               get_split_tile_size(tile,
-                                                                   max_render_feasible_tile_size);
-                                       vector<SplitRenderTile> to_path_trace_render_tiles =
-                                               split_tiles(tile, split_tile_size);
-                                       /* Print message to console */
-                                       if(background && (to_path_trace_render_tiles.size() > 1)) {
-                                               fprintf(stderr, "Message : Tiles need to be split "
-                                                       "further inside path trace (due to insufficient "
-                                                       "device-global-memory for split kernel to "
-                                                       "function) \n"
-                                                       "The current tile of dimensions %dx%d is split "
-                                                       "into tiles of dimension %dx%d for render \n",
-                                                       tile.w, tile.h,
-                                                       split_tile_size.x,
-                                                       split_tile_size.y);
-                                       }
-                                       /* Process all split tiles. */
-                                       for(int tile_iter = 0;
-                                           tile_iter < to_path_trace_render_tiles.size();
-                                           ++tile_iter)
-                                       {
-                                               path_trace(task,
-                                                          to_path_trace_render_tiles[tile_iter],
-                                                          max_render_feasible_tile_size);
-                                       }
-                               }
-                               else {
-                                       /* No splitting required; process the entire tile at once. */
-                                       /* Render feasible tile size is user-set-tile-size itself. */
-                                       max_render_feasible_tile_size.x =
-                                               (((tile_size.x - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
-                                               SPLIT_KERNEL_LOCAL_SIZE_X;
-                                       max_render_feasible_tile_size.y =
-                                               (((tile_size.y - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
-                                               SPLIT_KERNEL_LOCAL_SIZE_Y;
-                                       /* buffer_rng_state_stride is stride itself. */
-                                       SplitRenderTile split_tile(tile);
-                                       split_tile.buffer_rng_state_stride = tile.stride;
-                                       path_trace(task, split_tile, max_render_feasible_tile_size);
-                               }
-                               tile.sample = tile.start_sample + tile.num_samples;
-
-                               /* Complete kernel execution before release tile. */
-                               /* This helps in multi-device render;
-                                * The device that reaches the critical-section function
-                                * release_tile waits (stalling other devices from entering
-                                * release_tile) for all kernels to complete. If device1 (a
-                                * slow-render device) reaches release_tile first then it would
-                                * stall device2 (a fast-render device) from proceeding to render
-                                * next tile.
-                                */
-                               clFinish(cqCommandQueue);
-
-                               task->release_tile(tile);
-                       }
-               }
-       }
-
-protected:
-       cl_mem mem_alloc(size_t bufsize, cl_mem_flags mem_flag = CL_MEM_READ_WRITE)
-       {
-               cl_mem ptr;
-               assert(bufsize != 0);
-               ptr = clCreateBuffer(cxContext, mem_flag, bufsize, NULL, &ciErr);
-               opencl_assert_err(ciErr, "clCreateBuffer");
-               return ptr;
-       }
-
-       /* ** Those guys are for workign around some compiler-specific bugs ** */
-
-       cl_program load_cached_kernel(
-               const DeviceRequestedFeatures& /*requested_features*/,
-               OpenCLCache::ProgramName /*program_name*/,
-               thread_scoped_lock /*cache_locker*/)
-       {
-               VLOG(2) << "Skip loading kernel from cache, "
-                       << "not supported by split kernel.";
-               return NULL;
-       }
-
-       void store_cached_kernel(cl_platform_id /*platform*/,
-                                cl_device_id /*device*/,
-                                cl_program /*program*/,
-                                OpenCLCache::ProgramName /*program_name*/,
-                                thread_scoped_lock& /*slot_locker*/)
-       {
-               VLOG(2) << "Skip storing kernel in cache, "
-                       << "not supported by split kernel.";
-       }
-
-       string build_options_for_base_program(
-               const DeviceRequestedFeatures& requested_features)
-       {
-               return requested_features.get_build_options();
-       }
-};
-
 Device *device_opencl_create(DeviceInfo& info, Stats &stats, bool background)
 {
        vector<OpenCLPlatformDevice> usable_devices;
-       opencl_get_usable_devices(&usable_devices);
+       OpenCLInfo::get_usable_devices(&usable_devices);
        assert(info.num < usable_devices.size());
        const OpenCLPlatformDevice& platform_device = usable_devices[info.num];
        const string& platform_name = platform_device.platform_name;
        const cl_device_type device_type = platform_device.device_type;
-       if(opencl_kernel_use_split(platform_name, device_type)) {
+       if(OpenCLInfo::kernel_use_split(platform_name, device_type)) {
                VLOG(1) << "Using split kernel.";
-               return new OpenCLDeviceSplitKernel(info, stats, background);
+               return opencl_create_split_device(info, stats, background);
        } else {
                VLOG(1) << "Using mega kernel.";
-               return new OpenCLDeviceMegaKernel(info, stats, background);
+               return opencl_create_mega_device(info, stats, background);
        }
 }
 
@@ -3298,7 +52,7 @@ bool device_opencl_init(void)
 
        initialized = true;
 
-       if(opencl_device_type() != 0) {
+       if(OpenCLInfo::device_type() != 0) {
                int clew_result = clewInit();
                if(clew_result == CLEW_SUCCESS) {
                        VLOG(1) << "CLEW initialization succeeded.";
@@ -3322,7 +76,7 @@ bool device_opencl_init(void)
 void device_opencl_info(vector<DeviceInfo>& devices)
 {
        vector<OpenCLPlatformDevice> usable_devices;
-       opencl_get_usable_devices(&usable_devices);
+       OpenCLInfo::get_usable_devices(&usable_devices);
        /* Devices are numbered consecutively across platforms. */
        int num_devices = 0;
        foreach(OpenCLPlatformDevice& platform_device, usable_devices) {
@@ -3336,10 +90,10 @@ void device_opencl_info(vector<DeviceInfo>& devices)
                info.id = string_printf("OPENCL_%d", info.num);
                /* We don't know if it's used for display, but assume it is. */
                info.display_device = true;
-               info.advanced_shading = opencl_kernel_use_advanced_shading(platform_name);
+               info.advanced_shading = OpenCLInfo::kernel_use_advanced_shading(platform_name);
                info.pack_images = true;
-               info.use_split_kernel = opencl_kernel_use_split(platform_name,
-                                                               device_type);
+               info.use_split_kernel = OpenCLInfo::kernel_use_split(platform_name,
+                                                                    device_type);
                devices.push_back(info);
                num_devices++;
        }
@@ -3347,7 +101,7 @@ void device_opencl_info(vector<DeviceInfo>& devices)
 
 string device_opencl_capabilities(void)
 {
-       if(opencl_device_type() == 0) {
+       if(OpenCLInfo::device_type() == 0) {
                return "All OpenCL devices are forced to be OFF";
        }
        string result = "";
diff --git a/intern/cycles/device/opencl/opencl.h b/intern/cycles/device/opencl/opencl.h
new file mode 100644 (file)
index 0000000..032a97a
--- /dev/null
@@ -0,0 +1,397 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifdef WITH_OPENCL
+
+#include "clew.h"
+
+#include "device.h"
+
+#include "util_map.h"
+#include "util_param.h"
+#include "util_string.h"
+
+CCL_NAMESPACE_BEGIN
+
+#define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
+
+/* Macro declarations used with split kernel */
+
+/* Macro to enable/disable work-stealing */
+#define __WORK_STEALING__
+
+#define SPLIT_KERNEL_LOCAL_SIZE_X 64
+#define SPLIT_KERNEL_LOCAL_SIZE_Y 1
+
+/* This value may be tuned according to the scene we are rendering.
+ *
+ * Modifying PATH_ITER_INC_FACTOR value proportional to number of expected
+ * ray-bounces will improve performance.
+ */
+#define PATH_ITER_INC_FACTOR 8
+
+/* When allocate global memory in chunks. We may not be able to
+ * allocate exactly "CL_DEVICE_MAX_MEM_ALLOC_SIZE" bytes in chunks;
+ * Since some bytes may be needed for aligning chunks of memory;
+ * This is the amount of memory that we dedicate for that purpose.
+ */
+#define DATA_ALLOCATION_MEM_FACTOR 5000000 //5MB
+
+struct OpenCLPlatformDevice {
+       OpenCLPlatformDevice(cl_platform_id platform_id,
+                            const string& platform_name,
+                            cl_device_id device_id,
+                            cl_device_type device_type,
+                            const string& device_name)
+         : platform_id(platform_id),
+           platform_name(platform_name),
+           device_id(device_id),
+           device_type(device_type),
+           device_name(device_name) {}
+       cl_platform_id platform_id;
+       string platform_name;
+       cl_device_id device_id;
+       cl_device_type device_type;
+       string device_name;
+};
+
+/* Contains all static OpenCL helper functions. */
+class OpenCLInfo
+{
+public:
+       static cl_device_type device_type();
+       static bool use_debug();
+       static bool kernel_use_advanced_shading(const string& platform_name);
+       static bool kernel_use_split(const string& platform_name,
+                                    const cl_device_type device_type);
+       static bool device_supported(const string& platform_name,
+                                    const cl_device_id device_id);
+       static bool platform_version_check(cl_platform_id platform,
+                                          string *error = NULL);
+       static bool device_version_check(cl_device_id device,
+                                        string *error = NULL);
+       static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
+                                      bool force_all = false);
+};
+
+/* Thread safe cache for contexts and programs.
+ */
+class OpenCLCache
+{
+       struct Slot
+       {
+               struct ProgramEntry
+               {
+                       ProgramEntry();
+                       ProgramEntry(const ProgramEntry& rhs);
+                       ~ProgramEntry();
+                       cl_program program;
+                       thread_mutex *mutex;
+               };
+
+               Slot();
+               Slot(const Slot& rhs);
+               ~Slot();
+
+               thread_mutex *context_mutex;
+               cl_context context;
+               typedef map<ustring, ProgramEntry> EntryMap;
+               EntryMap programs;
+
+       };
+
+       /* key is combination of platform ID and device ID */
+       typedef pair<cl_platform_id, cl_device_id> PlatformDevicePair;
+
+       /* map of Slot objects */
+       typedef map<PlatformDevicePair, Slot> CacheMap;
+       CacheMap cache;
+
+       /* MD5 hash of the kernel source. */
+       string kernel_md5;
+
+       thread_mutex cache_lock;
+       thread_mutex kernel_md5_lock;
+
+       /* lazy instantiate */
+       static OpenCLCache& global_instance();
+
+public:
+
+       enum ProgramName {
+               OCL_DEV_BASE_PROGRAM,
+               OCL_DEV_MEGAKERNEL_PROGRAM,
+       };
+
+       /* Lookup context in the cache. If this returns NULL, slot_locker
+        * will be holding a lock for the cache. slot_locker should refer to a
+        * default constructed thread_scoped_lock. */
+       static cl_context get_context(cl_platform_id platform,
+                                     cl_device_id device,
+                                     thread_scoped_lock& slot_locker);
+       /* Same as above. */
+       static cl_program get_program(cl_platform_id platform,
+                                     cl_device_id device,
+                                     ustring key,
+                                     thread_scoped_lock& slot_locker);
+
+       /* Store context in the cache. You MUST have tried to get the item before storing to it. */
+       static void store_context(cl_platform_id platform,
+                                 cl_device_id device,
+                                 cl_context context,
+                                 thread_scoped_lock& slot_locker);
+       /* Same as above. */
+       static void store_program(cl_platform_id platform,
+                                 cl_device_id device,
+                                 cl_program program,
+                                 ustring key,
+                                 thread_scoped_lock& slot_locker);
+
+       static string get_kernel_md5();
+};
+
+#define opencl_assert(stmt) \
+       { \
+               cl_int err = stmt; \
+               \
+               if(err != CL_SUCCESS) { \
+                       string message = string_printf("OpenCL error: %s in %s", clewErrorString(err), #stmt); \
+                       if(error_msg == "") \
+                               error_msg = message; \
+                       fprintf(stderr, "%s\n", message.c_str()); \
+               } \
+       } (void)0
+
+class OpenCLDeviceBase : public Device
+{
+public:
+       DedicatedTaskPool task_pool;
+       cl_context cxContext;
+       cl_command_queue cqCommandQueue;
+       cl_platform_id cpPlatform;
+       cl_device_id cdDevice;
+       cl_int ciErr;
+
+       class OpenCLProgram {
+       public:
+               OpenCLProgram() : loaded(false), device(NULL) {}
+               OpenCLProgram(OpenCLDeviceBase *device,
+                             string program_name,
+                             string kernel_name,
+                             string kernel_build_options);
+               ~OpenCLProgram();
+
+               void add_kernel(ustring name);
+               void load();
+
+               bool is_loaded()    { return loaded; }
+               string get_log()    { return log; }
+               void report_error();
+
+               cl_kernel operator()();
+               cl_kernel operator()(ustring name);
+
+               void release();
+
+       private:
+               bool build_kernel(const string *debug_src);
+               bool compile_kernel(const string *debug_src);
+               bool load_binary(const string& clbin, const string *debug_src = NULL);
+               bool save_binary(const string& clbin);
+
+               bool loaded;
+               cl_program program;
+               OpenCLDeviceBase *device;
+
+               /* Used for the OpenCLCache key. */
+               string program_name;
+
+               string kernel_file, kernel_build_options, device_md5;
+               string error_msg, output_msg;
+               string log;
+
+               map<ustring, cl_kernel> kernels;
+       };
+
+       OpenCLProgram base_program;
+
+       typedef map<string, device_vector<uchar>*> ConstMemMap;
+       typedef map<string, device_ptr> MemMap;
+
+       ConstMemMap const_mem_map;
+       MemMap mem_map;
+       device_ptr null_mem;
+
+       bool device_initialized;
+       string platform_name;
+
+       bool opencl_error(cl_int err);
+       void opencl_error(const string& message);
+       void opencl_assert_err(cl_int err, const char* where);
+
+       OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_);
+       ~OpenCLDeviceBase();
+
+       static void CL_CALLBACK context_notify_callback(const char *err_info,
+               const void * /*private_info*/, size_t /*cb*/, void *user_data);
+
+       bool opencl_version_check();
+
+       string device_md5_hash(string kernel_custom_build_options = "");
+       bool load_kernels(const DeviceRequestedFeatures& requested_features);
+
+       /* Has to be implemented by the real device classes.
+        * The base device will then load all these programs. */
+       virtual void load_kernels(const DeviceRequestedFeatures& requested_features,
+                                 vector<OpenCLProgram*> &programs) = 0;
+
+       void mem_alloc(device_memory& mem, MemoryType type);
+       void mem_copy_to(device_memory& mem);
+       void mem_copy_from(device_memory& mem, int y, int w, int h, int elem);
+       void mem_zero(device_memory& mem);
+       void mem_free(device_memory& mem);
+       void const_copy_to(const char *name, void *host, size_t size);
+       void tex_alloc(const char *name,
+                      device_memory& mem,
+                      InterpolationType /*interpolation*/,
+                      ExtensionType /*extension*/);
+       void tex_free(device_memory& mem);
+
+       size_t global_size_round_up(int group_size, int global_size);
+       void enqueue_kernel(cl_kernel kernel, size_t w, size_t h);
+       void set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name);
+
+       void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half);
+       void shader(DeviceTask& task);
+
+       class OpenCLDeviceTask : public DeviceTask {
+       public:
+               OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
+               : DeviceTask(task)
+               {
+                       run = function_bind(&OpenCLDeviceBase::thread_run,
+                                           device,
+                                           this);
+               }
+       };
+
+       int get_split_task_count(DeviceTask& /*task*/)
+       {
+               return 1;
+       }
+
+       void task_add(DeviceTask& task)
+       {
+               task_pool.push(new OpenCLDeviceTask(this, task));
+       }
+
+       void task_wait()
+       {
+               task_pool.wait();
+       }
+
+       void task_cancel()
+       {
+               task_pool.cancel();
+       }
+
+       virtual void thread_run(DeviceTask * /*task*/) = 0;
+
+protected:
+       string kernel_build_options(const string *debug_src = NULL);
+
+       class ArgumentWrapper {
+       public:
+               ArgumentWrapper() : size(0), pointer(NULL) {}
+               template <typename T>
+               ArgumentWrapper(T& argument) : size(sizeof(argument)),
+                                              pointer(&argument) { }
+               ArgumentWrapper(int argument) : size(sizeof(int)),
+                                               int_value(argument),
+                                               pointer(&int_value) { }
+               ArgumentWrapper(float argument) : size(sizeof(float)),
+                                                 float_value(argument),
+                                                 pointer(&float_value) { }
+               size_t size;
+               int int_value;
+               float float_value;
+               void *pointer;
+       };
+
+       /* TODO(sergey): In the future we can use variadic templates, once
+        * C++0x is allowed. Should allow to clean this up a bit.
+        */
+       int kernel_set_args(cl_kernel kernel,
+                           int start_argument_index,
+                           const ArgumentWrapper& arg1 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg2 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg3 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg4 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg5 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg6 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg7 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg8 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg9 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg10 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg11 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg12 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg13 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg14 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg15 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg16 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg17 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg18 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg19 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg20 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg21 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg22 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg23 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg24 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg25 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg26 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg27 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg28 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg29 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg30 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg31 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg32 = ArgumentWrapper(),
+                           const ArgumentWrapper& arg33 = ArgumentWrapper());
+
+       void release_kernel_safe(cl_kernel kernel);
+       void release_mem_object_safe(cl_mem mem);
+       void release_program_safe(cl_program program);
+
+       /* ** Those guys are for workign around some compiler-specific bugs ** */
+
+       virtual cl_program load_cached_kernel(
+               ustring key,
+               thread_scoped_lock& cache_locker);
+
+       virtual void store_cached_kernel(
+               cl_program program,
+               ustring key,
+               thread_scoped_lock& cache_locker);
+
+       virtual string build_options_for_base_program(
+               const DeviceRequestedFeatures& /*requested_features*/);
+};
+
+Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, bool background);
+Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, bool background);
+
+CCL_NAMESPACE_END
+
+#endif
diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp
new file mode 100644 (file)
index 0000000..1e8fd26
--- /dev/null
@@ -0,0 +1,728 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifdef WITH_OPENCL
+
+#include "opencl.h"
+
+#include "kernel_types.h"
+
+#include "util_foreach.h"
+#include "util_logging.h"
+#include "util_md5.h"
+#include "util_path.h"
+#include "util_time.h"
+
+CCL_NAMESPACE_BEGIN
+
+bool OpenCLDeviceBase::opencl_error(cl_int err)
+{
+       if(err != CL_SUCCESS) {
+               string message = string_printf("OpenCL error (%d): %s", err, clewErrorString(err));
+               if(error_msg == "")
+                       error_msg = message;
+               fprintf(stderr, "%s\n", message.c_str());
+               return true;
+       }
+
+       return false;
+}
+
+void OpenCLDeviceBase::opencl_error(const string& message)
+{
+       if(error_msg == "")
+               error_msg = message;
+       fprintf(stderr, "%s\n", message.c_str());
+}
+
+void OpenCLDeviceBase::opencl_assert_err(cl_int err, const char* where)
+{
+       if(err != CL_SUCCESS) {
+               string message = string_printf("OpenCL error (%d): %s in %s", err, clewErrorString(err), where);
+               if(error_msg == "")
+                       error_msg = message;
+               fprintf(stderr, "%s\n", message.c_str());
+#ifndef NDEBUG
+               abort();
+#endif
+       }
+}
+
+OpenCLDeviceBase::OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
+: Device(info, stats, background_)
+{
+       cpPlatform = NULL;
+       cdDevice = NULL;
+       cxContext = NULL;
+       cqCommandQueue = NULL;
+       null_mem = 0;
+       device_initialized = false;
+
+       vector<OpenCLPlatformDevice> usable_devices;
+       OpenCLInfo::get_usable_devices(&usable_devices);
+       if(usable_devices.size() == 0) {
+               opencl_error("OpenCL: no devices found.");
+               return;
+       }
+       assert(info.num < usable_devices.size());
+       OpenCLPlatformDevice& platform_device = usable_devices[info.num];
+       cpPlatform = platform_device.platform_id;
+       cdDevice = platform_device.device_id;
+       platform_name = platform_device.platform_name;
+       VLOG(2) << "Creating new Cycles device for OpenCL platform "
+               << platform_name << ", device "
+               << platform_device.device_name << ".";
+
+       {
+               /* try to use cached context */
+               thread_scoped_lock cache_locker;
+               cxContext = OpenCLCache::get_context(cpPlatform, cdDevice, cache_locker);
+
+               if(cxContext == NULL) {
+                       /* create context properties array to specify platform */
+                       const cl_context_properties context_props[] = {
+                               CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform,
+                               0, 0
+                       };
+
+                       /* create context */
+                       cxContext = clCreateContext(context_props, 1, &cdDevice,
+                               context_notify_callback, cdDevice, &ciErr);
+
+                       if(opencl_error(ciErr)) {
+                               opencl_error("OpenCL: clCreateContext failed");
+                               return;
+                       }
+
+                       /* cache it */
+                       OpenCLCache::store_context(cpPlatform, cdDevice, cxContext, cache_locker);
+               }
+       }
+
+       cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
+       if(opencl_error(ciErr))
+               return;
+
+       null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
+       if(opencl_error(ciErr))
+               return;
+
+       fprintf(stderr, "Device init success\n");
+       device_initialized = true;
+}
+
+OpenCLDeviceBase::~OpenCLDeviceBase()
+{
+       task_pool.stop();
+
+       if(null_mem)
+               clReleaseMemObject(CL_MEM_PTR(null_mem));
+
+       ConstMemMap::iterator mt;
+       for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
+               mem_free(*(mt->second));
+               delete mt->second;
+       }
+
+       base_program.release();
+       if(cqCommandQueue)
+               clReleaseCommandQueue(cqCommandQueue);
+       if(cxContext)
+               clReleaseContext(cxContext);
+}
+
+void CL_CALLBACK OpenCLDeviceBase::context_notify_callback(const char *err_info,
+       const void * /*private_info*/, size_t /*cb*/, void *user_data)
+{
+       char name[256];
+       clGetDeviceInfo((cl_device_id)user_data, CL_DEVICE_NAME, sizeof(name), &name, NULL);
+
+       fprintf(stderr, "OpenCL error (%s): %s\n", name, err_info);
+}
+
+bool OpenCLDeviceBase::opencl_version_check()
+{
+       string error;
+       if(!OpenCLInfo::platform_version_check(cpPlatform, &error)) {
+               opencl_error(error);
+               return false;
+       }
+       if(!OpenCLInfo::device_version_check(cdDevice, &error)) {
+               opencl_error(error);
+               return false;
+       }
+       return true;
+}
+
+string OpenCLDeviceBase::device_md5_hash(string kernel_custom_build_options)
+{
+       MD5Hash md5;
+       char version[256], driver[256], name[256], vendor[256];
+
+       clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
+       clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
+       clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
+       clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
+
+       md5.append((uint8_t*)vendor, strlen(vendor));
+       md5.append((uint8_t*)version, strlen(version));
+       md5.append((uint8_t*)name, strlen(name));
+       md5.append((uint8_t*)driver, strlen(driver));
+
+       string options = kernel_build_options();
+       options += kernel_custom_build_options;
+       md5.append((uint8_t*)options.c_str(), options.size());
+
+       return md5.get_hex();
+}
+
+bool OpenCLDeviceBase::load_kernels(const DeviceRequestedFeatures& requested_features)
+{
+       /* Verify if device was initialized. */
+       if(!device_initialized) {
+               fprintf(stderr, "OpenCL: failed to initialize device.\n");
+               return false;
+       }
+
+       /* Verify we have right opencl version. */
+       if(!opencl_version_check())
+               return false;
+
+       base_program = OpenCLProgram(this, "base", "kernel.cl", build_options_for_base_program(requested_features));
+       base_program.add_kernel(ustring("convert_to_byte"));
+       base_program.add_kernel(ustring("convert_to_half_float"));
+       base_program.add_kernel(ustring("shader"));
+       base_program.add_kernel(ustring("bake"));
+
+       vector<OpenCLProgram*> programs;
+       programs.push_back(&base_program);
+       /* Call actual class to fill the vector with its programs. */
+       load_kernels(requested_features, programs);
+
+       TaskPool task_pool;
+       foreach(OpenCLProgram *program, programs) {
+               task_pool.push(function_bind(&OpenCLProgram::load, program));
+       }
+       task_pool.wait_work();
+
+       foreach(OpenCLProgram *program, programs) {
+               VLOG(2) << program->get_log();
+               if(!program->is_loaded()) {
+                       program->report_error();
+                       return false;
+               }
+       }
+
+       return true;
+}
+
+void OpenCLDeviceBase::mem_alloc(device_memory& mem, MemoryType type)
+{
+       size_t size = mem.memory_size();
+
+       cl_mem_flags mem_flag;
+       void *mem_ptr = NULL;
+
+       if(type == MEM_READ_ONLY)
+               mem_flag = CL_MEM_READ_ONLY;
+       else if(type == MEM_WRITE_ONLY)
+               mem_flag = CL_MEM_WRITE_ONLY;
+       else
+               mem_flag = CL_MEM_READ_WRITE;
+
+       /* Zero-size allocation might be invoked by render, but not really
+        * supported by OpenCL. Using NULL as device pointer also doesn't really
+        * work for some reason, so for the time being we'll use special case
+        * will null_mem buffer.
+        */
+       if(size != 0) {
+               mem.device_pointer = (device_ptr)clCreateBuffer(cxContext,
+                                                               mem_flag,
+                                                               size,
+                                                               mem_ptr,
+                                                               &ciErr);
+               opencl_assert_err(ciErr, "clCreateBuffer");
+       }
+       else {
+               mem.device_pointer = null_mem;
+       }
+
+       stats.mem_alloc(size);
+       mem.device_size = size;
+}
+
+void OpenCLDeviceBase::mem_copy_to(device_memory& mem)
+{
+       /* this is blocking */
+       size_t size = mem.memory_size();
+       if(size != 0) {
+               opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
+                                                  CL_MEM_PTR(mem.device_pointer),
+                                                  CL_TRUE,
+                                                  0,
+                                                  size,
+                                                  (void*)mem.data_pointer,
+                                                  0,
+                                                  NULL, NULL));
+       }
+}
+
+void OpenCLDeviceBase::mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
+{
+       size_t offset = elem*y*w;
+       size_t size = elem*w*h;
+       assert(size != 0);
+       opencl_assert(clEnqueueReadBuffer(cqCommandQueue,
+                                         CL_MEM_PTR(mem.device_pointer),
+                                         CL_TRUE,
+                                         offset,
+                                         size,
+                                         (uchar*)mem.data_pointer + offset,
+                                         0,
+                                         NULL, NULL));
+}
+
+void OpenCLDeviceBase::mem_zero(device_memory& mem)
+{
+       if(mem.device_pointer) {
+               memset((void*)mem.data_pointer, 0, mem.memory_size());
+               mem_copy_to(mem);
+       }
+}
+
+void OpenCLDeviceBase::mem_free(device_memory& mem)
+{
+       if(mem.device_pointer) {
+               if(mem.device_pointer != null_mem) {
+                       opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
+               }
+               mem.device_pointer = 0;
+
+               stats.mem_free(mem.device_size);
+               mem.device_size = 0;
+       }
+}
+
+void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
+{
+       ConstMemMap::iterator i = const_mem_map.find(name);
+
+       if(i == const_mem_map.end()) {
+               device_vector<uchar> *data = new device_vector<uchar>();
+               data->copy((uchar*)host, size);
+
+               mem_alloc(*data, MEM_READ_ONLY);
+               i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first;
+       }
+       else {
+               device_vector<uchar> *data = i->second;
+               data->copy((uchar*)host, size);
+       }
+
+       mem_copy_to(*i->second);
+}
+
+void OpenCLDeviceBase::tex_alloc(const char *name,
+               device_memory& mem,
+               InterpolationType /*interpolation*/,
+               ExtensionType /*extension*/)
+{
+       VLOG(1) << "Texture allocate: " << name << ", "
+               << string_human_readable_number(mem.memory_size()) << " bytes. ("
+               << string_human_readable_size(mem.memory_size()) << ")";
+       mem_alloc(mem, MEM_READ_ONLY);
+       mem_copy_to(mem);
+       assert(mem_map.find(name) == mem_map.end());
+       mem_map.insert(MemMap::value_type(name, mem.device_pointer));
+}
+
+void OpenCLDeviceBase::tex_free(device_memory& mem)
+{
+       if(mem.device_pointer) {
+               foreach(const MemMap::value_type& value, mem_map) {
+                       if(value.second == mem.device_pointer) {
+                               mem_map.erase(value.first);
+                               break;
+                       }
+               }
+
+               mem_free(mem);
+       }
+}
+
+size_t OpenCLDeviceBase::global_size_round_up(int group_size, int global_size)
+{
+       int r = global_size % group_size;
+       return global_size + ((r == 0)? 0: group_size - r);
+}
+
+void OpenCLDeviceBase::enqueue_kernel(cl_kernel kernel, size_t w, size_t h)
+{
+       size_t workgroup_size, max_work_items[3];
+
+       clGetKernelWorkGroupInfo(kernel, cdDevice,
+               CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
+       clGetDeviceInfo(cdDevice,
+               CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, max_work_items, NULL);
+
+       /* Try to divide evenly over 2 dimensions. */
+       size_t sqrt_workgroup_size = max((size_t)sqrt((double)workgroup_size), 1);
+       size_t local_size[2] = {sqrt_workgroup_size, sqrt_workgroup_size};
+
+       /* Some implementations have max size 1 on 2nd dimension. */
+       if(local_size[1] > max_work_items[1]) {
+               local_size[0] = workgroup_size/max_work_items[1];
+               local_size[1] = max_work_items[1];
+       }
+
+       size_t global_size[2] = {global_size_round_up(local_size[0], w),
+                                global_size_round_up(local_size[1], h)};
+
+       /* Vertical size of 1 is coming from bake/shade kernels where we should
+        * not round anything up because otherwise we'll either be doing too
+        * much work per pixel (if we don't check global ID on Y axis) or will
+        * be checking for global ID to always have Y of 0.
+        */
+       if (h == 1) {
+               global_size[h] = 1;
+       }
+
+       /* run kernel */
+       opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));
+       opencl_assert(clFlush(cqCommandQueue));
+}
+
+void OpenCLDeviceBase::set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name)
+{
+       cl_mem ptr;
+
+       MemMap::iterator i = mem_map.find(name);
+       if(i != mem_map.end()) {
+               ptr = CL_MEM_PTR(i->second);
+       }
+       else {
+               /* work around NULL not working, even though the spec says otherwise */
+               ptr = CL_MEM_PTR(null_mem);
+       }
+
+       opencl_assert(clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr));
+}
+
+void OpenCLDeviceBase::film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
+{
+       /* cast arguments to cl types */
+       cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+       cl_mem d_rgba = (rgba_byte)? CL_MEM_PTR(rgba_byte): CL_MEM_PTR(rgba_half);
+       cl_mem d_buffer = CL_MEM_PTR(buffer);
+       cl_int d_x = task.x;
+       cl_int d_y = task.y;
+       cl_int d_w = task.w;
+       cl_int d_h = task.h;
+       cl_float d_sample_scale = 1.0f/(task.sample + 1);
+       cl_int d_offset = task.offset;
+       cl_int d_stride = task.stride;
+
+
+       cl_kernel ckFilmConvertKernel = (rgba_byte)? base_program(ustring("convert_to_byte")): base_program(ustring("convert_to_half_float"));
+
+       cl_uint start_arg_index =
+               kernel_set_args(ckFilmConvertKernel,
+                               0,
+                               d_data,
+                               d_rgba,
+                               d_buffer);
+
+#define KERNEL_TEX(type, ttype, name) \
+set_kernel_arg_mem(ckFilmConvertKernel, &start_arg_index, #name);
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+
+       start_arg_index += kernel_set_args(ckFilmConvertKernel,
+                                          start_arg_index,
+                                          d_sample_scale,
+                                          d_x,
+                                          d_y,
+                                          d_w,
+                                          d_h,
+                                          d_offset,
+                                          d_stride);
+
+       enqueue_kernel(ckFilmConvertKernel, d_w, d_h);
+}
+
+void OpenCLDeviceBase::shader(DeviceTask& task)
+{
+       /* cast arguments to cl types */
+       cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+       cl_mem d_input = CL_MEM_PTR(task.shader_input);
+       cl_mem d_output = CL_MEM_PTR(task.shader_output);
+       cl_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma);
+       cl_int d_shader_eval_type = task.shader_eval_type;
+       cl_int d_shader_filter = task.shader_filter;
+       cl_int d_shader_x = task.shader_x;
+       cl_int d_shader_w = task.shader_w;
+       cl_int d_offset = task.offset;
+
+       cl_kernel kernel;
+
+       if(task.shader_eval_type >= SHADER_EVAL_BAKE)
+               kernel = base_program(ustring("shader"));
+       else
+               kernel = base_program(ustring("bake"));
+
+       cl_uint start_arg_index =
+               kernel_set_args(kernel,
+                               0,
+                               d_data,
+                               d_input,
+                               d_output);
+
+       if(task.shader_eval_type < SHADER_EVAL_BAKE) {
+               start_arg_index += kernel_set_args(kernel,
+                                                  start_arg_index,
+                                                  d_output_luma);
+       }
+
+#define KERNEL_TEX(type, ttype, name) \
+       set_kernel_arg_mem(kernel, &start_arg_index, #name);
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+
+       start_arg_index += kernel_set_args(kernel,
+                                          start_arg_index,
+                                          d_shader_eval_type);
+       if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
+               start_arg_index += kernel_set_args(kernel,
+                                                  start_arg_index,
+                                                  d_shader_filter);
+       }
+       start_arg_index += kernel_set_args(kernel,
+                                          start_arg_index,
+                                          d_shader_x,
+                                          d_shader_w,
+                                          d_offset);
+
+       for(int sample = 0; sample < task.num_samples; sample++) {
+
+               if(task.get_cancel())
+                       break;
+
+               kernel_set_args(kernel, start_arg_index, sample);
+
+               enqueue_kernel(kernel, task.shader_w, 1);
+
+               clFinish(cqCommandQueue);
+
+               task.update_progress(NULL);
+       }
+}
+
+string OpenCLDeviceBase::kernel_build_options(const string *debug_src)
+{
+       string build_options = "-cl-fast-relaxed-math ";
+
+       if(platform_name == "NVIDIA CUDA") {
+               build_options += "-D__KERNEL_OPENCL_NVIDIA__ "
+                                "-cl-nv-maxrregcount=32 "
+                                "-cl-nv-verbose ";
+
+               uint compute_capability_major, compute_capability_minor;
+               clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,
+                               sizeof(cl_uint), &compute_capability_major, NULL);
+               clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV,
+                               sizeof(cl_uint), &compute_capability_minor, NULL);
+
+               build_options += string_printf("-D__COMPUTE_CAPABILITY__=%u ",
+                                              compute_capability_major * 100 +
+                                              compute_capability_minor * 10);
+       }
+
+       else if(platform_name == "Apple")
+               build_options += "-D__KERNEL_OPENCL_APPLE__ ";
+
+       else if(platform_name == "AMD Accelerated Parallel Processing")
+               build_options += "-D__KERNEL_OPENCL_AMD__ ";
+
+       else if(platform_name == "Intel(R) OpenCL") {
+               build_options += "-D__KERNEL_OPENCL_INTEL_CPU__ ";
+
+               /* Options for gdb source level kernel debugging.
+                * this segfaults on linux currently.
+                */
+               if(OpenCLInfo::use_debug() && debug_src)
+                       build_options += "-g -s \"" + *debug_src + "\" ";
+       }
+
+       if(OpenCLInfo::use_debug())
+               build_options += "-D__KERNEL_OPENCL_DEBUG__ ";
+
+#ifdef WITH_CYCLES_DEBUG
+       build_options += "-D__KERNEL_DEBUG__ ";
+#endif
+
+       return build_options;
+}
+
+/* TODO(sergey): In the future we can use variadic templates, once
+ * C++0x is allowed. Should allow to clean this up a bit.
+ */
+int OpenCLDeviceBase::kernel_set_args(cl_kernel kernel,
+                    int start_argument_index,
+                    const ArgumentWrapper& arg1,
+                    const ArgumentWrapper& arg2,
+                    const ArgumentWrapper& arg3,
+                    const ArgumentWrapper& arg4,
+                    const ArgumentWrapper& arg5,
+                    const ArgumentWrapper& arg6,
+                    const ArgumentWrapper& arg7,
+                    const ArgumentWrapper& arg8,
+                    const ArgumentWrapper& arg9,
+                    const ArgumentWrapper& arg10,
+                    const ArgumentWrapper& arg11,
+                    const ArgumentWrapper& arg12,
+                    const ArgumentWrapper& arg13,
+                    const ArgumentWrapper& arg14,
+                    const ArgumentWrapper& arg15,
+                    const ArgumentWrapper& arg16,
+                    const ArgumentWrapper& arg17,
+                    const ArgumentWrapper& arg18,
+                    const ArgumentWrapper& arg19,
+                    const ArgumentWrapper& arg20,
+                    const ArgumentWrapper& arg21,
+                    const ArgumentWrapper& arg22,
+                    const ArgumentWrapper& arg23,
+                    const ArgumentWrapper& arg24,
+                    const ArgumentWrapper& arg25,
+                    const ArgumentWrapper& arg26,
+                    const ArgumentWrapper& arg27,
+                    const ArgumentWrapper& arg28,
+                    const ArgumentWrapper& arg29,
+                    const ArgumentWrapper& arg30,
+                    const ArgumentWrapper& arg31,
+                    const ArgumentWrapper& arg32,
+                    const ArgumentWrapper& arg33)
+{
+       int current_arg_index = 0;
+#define FAKE_VARARG_HANDLE_ARG(arg) \
+       do { \
+               if(arg.pointer != NULL) { \
+                       opencl_assert(clSetKernelArg( \
+                               kernel, \
+                               start_argument_index + current_arg_index, \
+                               arg.size, arg.pointer)); \
+                       ++current_arg_index; \
+               } \
+               else { \
+                       return current_arg_index; \
+               } \
+       } while(false)
+       FAKE_VARARG_HANDLE_ARG(arg1);
+       FAKE_VARARG_HANDLE_ARG(arg2);
+       FAKE_VARARG_HANDLE_ARG(arg3);
+       FAKE_VARARG_HANDLE_ARG(arg4);
+       FAKE_VARARG_HANDLE_ARG(arg5);
+       FAKE_VARARG_HANDLE_ARG(arg6);
+       FAKE_VARARG_HANDLE_ARG(arg7);
+       FAKE_VARARG_HANDLE_ARG(arg8);
+       FAKE_VARARG_HANDLE_ARG(arg9);
+       FAKE_VARARG_HANDLE_ARG(arg10);
+       FAKE_VARARG_HANDLE_ARG(arg11);
+       FAKE_VARARG_HANDLE_ARG(arg12);
+       FAKE_VARARG_HANDLE_ARG(arg13);
+       FAKE_VARARG_HANDLE_ARG(arg14);
+       FAKE_VARARG_HANDLE_ARG(arg15);
+       FAKE_VARARG_HANDLE_ARG(arg16);
+       FAKE_VARARG_HANDLE_ARG(arg17);
+       FAKE_VARARG_HANDLE_ARG(arg18);
+       FAKE_VARARG_HANDLE_ARG(arg19);
+       FAKE_VARARG_HANDLE_ARG(arg20);
+       FAKE_VARARG_HANDLE_ARG(arg21);
+       FAKE_VARARG_HANDLE_ARG(arg22);
+       FAKE_VARARG_HANDLE_ARG(arg23);
+       FAKE_VARARG_HANDLE_ARG(arg24);
+       FAKE_VARARG_HANDLE_ARG(arg25);
+       FAKE_VARARG_HANDLE_ARG(arg26);
+       FAKE_VARARG_HANDLE_ARG(arg27);
+       FAKE_VARARG_HANDLE_ARG(arg28);
+       FAKE_VARARG_HANDLE_ARG(arg29);
+       FAKE_VARARG_HANDLE_ARG(arg30);
+       FAKE_VARARG_HANDLE_ARG(arg31);
+       FAKE_VARARG_HANDLE_ARG(arg32);
+       FAKE_VARARG_HANDLE_ARG(arg33);
+#undef FAKE_VARARG_HANDLE_ARG
+       return current_arg_index;
+}
+
+void OpenCLDeviceBase::release_kernel_safe(cl_kernel kernel)
+{
+       if(kernel) {
+               clReleaseKernel(kernel);
+       }
+}
+
+void OpenCLDeviceBase::release_mem_object_safe(cl_mem mem)
+{
+       if(mem != NULL) {
+               clReleaseMemObject(mem);
+       }
+}
+
+void OpenCLDeviceBase::release_program_safe(cl_program program)
+{
+       if(program) {
+               clReleaseProgram(program);
+       }
+}
+
+/* ** Those guys are for workign around some compiler-specific bugs ** */
+
+cl_program OpenCLDeviceBase::load_cached_kernel(
+        ustring key,
+        thread_scoped_lock& cache_locker)
+{
+       return OpenCLCache::get_program(cpPlatform,
+                                       cdDevice,
+                                       key,
+                                       cache_locker);
+}
+
+void OpenCLDeviceBase::store_cached_kernel(
+        cl_program program,
+        ustring key,
+        thread_scoped_lock& cache_locker)
+{
+       OpenCLCache::store_program(cpPlatform,
+                                  cdDevice,
+                                  program,
+                                  key,
+                                  cache_locker);
+}
+
+string OpenCLDeviceBase::build_options_for_base_program(
+        const DeviceRequestedFeatures& /*requested_features*/)
+{
+       /* TODO(sergey): By default we compile all features, meaning
+        * mega kernel is not getting feature-based optimizations.
+        *
+        * Ideally we need always compile kernel with as less features
+        * enabled as possible to keep performance at it's max.
+        */
+       return "";
+}
+
+CCL_NAMESPACE_END
+
+#endif
diff --git a/intern/cycles/device/opencl/opencl_mega.cpp b/intern/cycles/device/opencl/opencl_mega.cpp
new file mode 100644 (file)
index 0000000..369c086
--- /dev/null
@@ -0,0 +1,150 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifdef WITH_OPENCL
+
+#include "opencl.h"
+
+#include "buffers.h"
+
+#include "kernel_types.h"
+
+#include "util_md5.h"
+#include "util_path.h"
+#include "util_time.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OpenCLDeviceMegaKernel : public OpenCLDeviceBase
+{
+public:
+       OpenCLProgram path_trace_program;
+
+       OpenCLDeviceMegaKernel(DeviceInfo& info, Stats &stats, bool background_)
+       : OpenCLDeviceBase(info, stats, background_),
+         path_trace_program(this, "megakernel", "kernel.cl", "-D__COMPILE_ONLY_MEGAKERNEL__ ")
+       {
+       }
+
+       virtual void load_kernels(const DeviceRequestedFeatures& /*requested_features*/,
+                                 vector<OpenCLProgram*> &programs)
+       {
+               path_trace_program.add_kernel(ustring("path_trace"));
+               programs.push_back(&path_trace_program);
+       }
+
+       ~OpenCLDeviceMegaKernel()
+       {
+               task_pool.stop();
+               path_trace_program.release();
+       }
+
+       void path_trace(RenderTile& rtile, int sample)
+       {
+               /* Cast arguments to cl types. */
+               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+               cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
+               cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
+               cl_int d_x = rtile.x;
+               cl_int d_y = rtile.y;
+               cl_int d_w = rtile.w;
+               cl_int d_h = rtile.h;
+               cl_int d_offset = rtile.offset;
+               cl_int d_stride = rtile.stride;
+
+               /* Sample arguments. */
+               cl_int d_sample = sample;
+
+               cl_kernel ckPathTraceKernel = path_trace_program(ustring("path_trace"));
+
+               cl_uint start_arg_index =
+                       kernel_set_args(ckPathTraceKernel,
+                                       0,
+                                       d_data,
+                                       d_buffer,
+                                       d_rng_state);
+
+#define KERNEL_TEX(type, ttype, name) \
+               set_kernel_arg_mem(ckPathTraceKernel, &start_arg_index, #name);
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+
+               start_arg_index += kernel_set_args(ckPathTraceKernel,
+                                                  start_arg_index,
+                                                  d_sample,
+                                                  d_x,
+                                                  d_y,
+                                                  d_w,
+                                                  d_h,
+                                                  d_offset,
+                                                  d_stride);
+
+               enqueue_kernel(ckPathTraceKernel, d_w, d_h);
+       }
+
+       void thread_run(DeviceTask *task)
+       {
+               if(task->type == DeviceTask::FILM_CONVERT) {
+                       film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
+               }
+               else if(task->type == DeviceTask::SHADER) {
+                       shader(*task);
+               }
+               else if(task->type == DeviceTask::PATH_TRACE) {
+                       RenderTile tile;
+                       /* Keep rendering tiles until done. */
+                       while(task->acquire_tile(this, tile)) {
+                               int start_sample = tile.start_sample;
+                               int end_sample = tile.start_sample + tile.num_samples;
+
+                               for(int sample = start_sample; sample < end_sample; sample++) {
+                                       if(task->get_cancel()) {
+                                               if(task->need_finish_queue == false)
+                                                       break;
+                                       }
+
+                                       path_trace(tile, sample);
+
+                                       tile.sample = sample + 1;
+
+                                       task->update_progress(&tile);
+                               }
+
+                               /* Complete kernel execution before release tile */
+                               /* This helps in multi-device render;
+                                * The device that reaches the critical-section function
+                                * release_tile waits (stalling other devices from entering
+                                * release_tile) for all kernels to complete. If device1 (a
+                                * slow-render device) reaches release_tile first then it would
+                                * stall device2 (a fast-render device) from proceeding to render
+                                * next tile.
+                                */
+                               clFinish(cqCommandQueue);
+
+                               task->release_tile(tile);
+                       }
+               }
+       }
+};
+
+Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, bool background)
+{
+       return new OpenCLDeviceMegaKernel(info, stats, background);
+}
+
+CCL_NAMESPACE_END
+
+#endif
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
new file mode 100644 (file)
index 0000000..239e73a
--- /dev/null
@@ -0,0 +1,1307 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifdef WITH_OPENCL
+
+#include "opencl.h"
+
+#include "buffers.h"
+
+#include "kernel_types.h"
+
+#include "util_md5.h"
+#include "util_path.h"
+#include "util_time.h"
+
+CCL_NAMESPACE_BEGIN
+
+/* TODO(sergey): This is to keep tile split on OpenCL level working
+ * for now, since without this view-port render does not work as it
+ * should.
+ *
+ * Ideally it'll be done on the higher level, but we need to get ready
+ * for merge rather soon, so let's keep split logic private here in
+ * the file.
+ */
+class SplitRenderTile : public RenderTile {
+public:
+       SplitRenderTile()
+               : RenderTile(),
+                 buffer_offset_x(0),
+                 buffer_offset_y(0),
+                 rng_state_offset_x(0),
+                 rng_state_offset_y(0),
+                 buffer_rng_state_stride(0) {}
+
+       explicit SplitRenderTile(RenderTile& tile)
+               : RenderTile(),
+                 buffer_offset_x(0),
+                 buffer_offset_y(0),
+                 rng_state_offset_x(0),
+                 rng_state_offset_y(0),
+                 buffer_rng_state_stride(0)
+       {
+               x = tile.x;
+               y = tile.y;
+               w = tile.w;
+               h = tile.h;
+               start_sample = tile.start_sample;
+               num_samples = tile.num_samples;
+               sample = tile.sample;
+               resolution = tile.resolution;
+               offset = tile.offset;
+               stride = tile.stride;
+               buffer = tile.buffer;
+               rng_state = tile.rng_state;
+               buffers = tile.buffers;
+       }
+
+       /* Split kernel is device global memory constrained;
+        * hence split kernel cant render big tile size's in
+        * one go. If the user sets a big tile size (big tile size
+        * is a term relative to the available device global memory),
+        * we split the tile further and then call path_trace on
+        * each of those split tiles. The following variables declared,
+        * assist in achieving that purpose
+        */
+       int buffer_offset_x;
+       int buffer_offset_y;
+       int rng_state_offset_x;
+       int rng_state_offset_y;
+       int buffer_rng_state_stride;
+};
+
+/* OpenCLDeviceSplitKernel's declaration/definition. */
+class OpenCLDeviceSplitKernel : public OpenCLDeviceBase
+{
+public:
+       /* Kernel declaration. */
+       OpenCLProgram program_data_init;
+       OpenCLProgram program_scene_intersect;
+       OpenCLProgram program_lamp_emission;
+       OpenCLProgram program_queue_enqueue;
+       OpenCLProgram program_background_buffer_update;
+       OpenCLProgram program_shader_eval;
+       OpenCLProgram program_holdout_emission_blurring_pathtermination_ao;
+       OpenCLProgram program_direct_lighting;
+       OpenCLProgram program_shadow_blocked;
+       OpenCLProgram program_next_iteration_setup;
+       OpenCLProgram program_sum_all_radiance;
+
+       /* Global memory variables [porting]; These memory is used for
+        * co-operation between different kernels; Data written by one
+        * kernel will be available to another kernel via this global
+        * memory.
+        */
+       cl_mem rng_coop;
+       cl_mem throughput_coop;
+       cl_mem L_transparent_coop;
+       cl_mem PathRadiance_coop;
+       cl_mem Ray_coop;
+       cl_mem PathState_coop;
+       cl_mem Intersection_coop;
+       cl_mem kgbuffer;  /* KernelGlobals buffer. */
+
+       /* Global buffers for ShaderData. */
+       cl_mem sd;             /* ShaderData used in the main path-iteration loop. */
+       cl_mem sd_DL_shadow;   /* ShaderData used in Direct Lighting and
+                               * shadow_blocked kernel.
+                               */
+
+       /* Global memory required for shadow blocked and accum_radiance. */
+       cl_mem BSDFEval_coop;
+       cl_mem ISLamp_coop;
+       cl_mem LightRay_coop;
+       cl_mem AOAlpha_coop;
+       cl_mem AOBSDF_coop;
+       cl_mem AOLightRay_coop;
+       cl_mem Intersection_coop_shadow;
+
+#ifdef WITH_CYCLES_DEBUG
+       /* DebugData memory */
+       cl_mem debugdata_coop;
+#endif
+
+       /* Global state array that tracks ray state. */
+       cl_mem ray_state;
+
+       /* Per sample buffers. */
+       cl_mem per_sample_output_buffers;
+
+       /* Denotes which sample each ray is being processed for. */
+       cl_mem work_array;
+
+       /* Queue */
+       cl_mem Queue_data;  /* Array of size queuesize * num_queues * sizeof(int). */
+       cl_mem Queue_index; /* Array of size num_queues * sizeof(int);
+                            * Tracks the size of each queue.
+                            */
+
+       /* Flag to make sceneintersect and lampemission kernel use queues. */
+       cl_mem use_queues_flag;
+
+       /* Amount of memory in output buffer associated with one pixel/thread. */
+       size_t per_thread_output_buffer_size;
+
+       /* Total allocatable available device memory. */
+       size_t total_allocatable_memory;
+
+       /* host version of ray_state; Used in checking host path-iteration
+        * termination.
+        */
+       char *hostRayStateArray;
+
+       /* Number of path-iterations to be done in one shot. */
+       unsigned int PathIteration_times;
+
+#ifdef __WORK_STEALING__
+       /* Work pool with respect to each work group. */
+       cl_mem work_pool_wgs;
+
+       /* Denotes the maximum work groups possible w.r.t. current tile size. */
+       unsigned int max_work_groups;
+#endif
+
+       /* clos_max value for which the kernels have been loaded currently. */
+       int current_max_closure;
+
+       /* Marked True in constructor and marked false at the end of path_trace(). */
+       bool first_tile;
+
+       OpenCLDeviceSplitKernel(DeviceInfo& info, Stats &stats, bool background_)
+       : OpenCLDeviceBase(info, stats, background_)
+       {
+               background = background_;
+
+               /* Initialize cl_mem variables. */
+               kgbuffer = NULL;
+               sd = NULL;
+               sd_DL_shadow = NULL;
+
+               rng_coop = NULL;
+               throughput_coop = NULL;
+               L_transparent_coop = NULL;
+               PathRadiance_coop = NULL;
+               Ray_coop = NULL;
+               PathState_coop = NULL;
+               Intersection_coop = NULL;
+               ray_state = NULL;
+
+               AOAlpha_coop = NULL;
+               AOBSDF_coop = NULL;
+               AOLightRay_coop = NULL;
+               BSDFEval_coop = NULL;
+               ISLamp_coop = NULL;
+               LightRay_coop = NULL;
+               Intersection_coop_shadow = NULL;
+
+#ifdef WITH_CYCLES_DEBUG
+               debugdata_coop = NULL;
+#endif
+
+               work_array = NULL;
+
+               /* Queue. */
+               Queue_data = NULL;
+               Queue_index = NULL;
+               use_queues_flag = NULL;
+
+               per_sample_output_buffers = NULL;
+
+               per_thread_output_buffer_size = 0;
+               hostRayStateArray = NULL;
+               PathIteration_times = PATH_ITER_INC_FACTOR;
+#ifdef __WORK_STEALING__
+               work_pool_wgs = NULL;
+               max_work_groups = 0;
+#endif
+               current_max_closure = -1;
+               first_tile = true;
+
+               /* Get device's maximum memory that can be allocated. */
+               ciErr = clGetDeviceInfo(cdDevice,
+                                       CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                                       sizeof(size_t),
+                                       &total_allocatable_memory,
+                                       NULL);
+               assert(ciErr == CL_SUCCESS);
+               if(platform_name == "AMD Accelerated Parallel Processing") {
+                       /* This value is tweak-able; AMD platform does not seem to
+                        * give maximum performance when all of CL_DEVICE_MAX_MEM_ALLOC_SIZE
+                        * is considered for further computation.
+                        */
+                       total_allocatable_memory /= 2;
+               }
+       }
+
+       /* Split kernel utility functions. */
+       size_t get_tex_size(const char *tex_name)
+       {
+               cl_mem ptr;
+               size_t ret_size = 0;
+               MemMap::iterator i = mem_map.find(tex_name);
+               if(i != mem_map.end()) {
+                       ptr = CL_MEM_PTR(i->second);
+                       ciErr = clGetMemObjectInfo(ptr,
+                                                  CL_MEM_SIZE,
+                                                  sizeof(ret_size),
+                                                  &ret_size,
+                                                  NULL);
+                       assert(ciErr == CL_SUCCESS);
+               }
+               return ret_size;
+       }
+
+       size_t get_shader_data_size(size_t max_closure)
+       {
+               /* ShaderData size with variable size ShaderClosure array */
+               return sizeof(ShaderData) - (sizeof(ShaderClosure) * (MAX_CLOSURE - max_closure));
+       }
+
+       /* Returns size of KernelGlobals structure associated with OpenCL. */
+       size_t get_KernelGlobals_size()
+       {
+               /* Copy dummy KernelGlobals related to OpenCL from kernel_globals.h to
+                * fetch its size.
+                */
+               typedef struct KernelGlobals {
+                       ccl_constant KernelData *data;
+#define KERNEL_TEX(type, ttype, name) \
+       ccl_global type *name;
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+                       void *sd_input;
+                       void *isect_shadow;
+               } KernelGlobals;
+
+               return sizeof(KernelGlobals);
+       }
+
+       virtual void load_kernels(const DeviceRequestedFeatures& requested_features,
+                                 vector<OpenCLProgram*> &programs)
+       {
+               string build_options = "-D__SPLIT_KERNEL__ ";
+#ifdef __WORK_STEALING__
+               build_options += "-D__WORK_STEALING__ ";
+#endif
+               build_options += requested_features.get_build_options();
+
+               /* Set compute device build option. */
+               cl_device_type device_type;
+               ciErr = clGetDeviceInfo(cdDevice,
+                                       CL_DEVICE_TYPE,
+                                       sizeof(cl_device_type),
+                                       &device_type,
+                                       NULL);
+               assert(ciErr == CL_SUCCESS);
+               if(device_type == CL_DEVICE_TYPE_GPU) {
+                       build_options += " -D__COMPUTE_DEVICE_GPU__";
+               }
+
+#define GLUE(a, b) a ## b
+#define LOAD_KERNEL(name) \
+       do { \
+               GLUE(program_, name) = OpenCLProgram(this, "split_" #name, "kernel_" #name ".cl", build_options); \
+               GLUE(program_, name).add_kernel(ustring("path_trace_" #name)); \
+               programs.push_back(&GLUE(program_, name)); \
+       } while(false)
+
+               LOAD_KERNEL(data_init);
+               LOAD_KERNEL(scene_intersect);
+               LOAD_KERNEL(lamp_emission);
+               LOAD_KERNEL(queue_enqueue);
+               LOAD_KERNEL(background_buffer_update);
+               LOAD_KERNEL(shader_eval);
+               LOAD_KERNEL(holdout_emission_blurring_pathtermination_ao);
+               LOAD_KERNEL(direct_lighting);
+               LOAD_KERNEL(shadow_blocked);
+               LOAD_KERNEL(next_iteration_setup);
+               LOAD_KERNEL(sum_all_radiance);
+
+#undef FIND_KERNEL
+#undef GLUE
+
+               current_max_closure = requested_features.max_closure;
+       }
+
+       ~OpenCLDeviceSplitKernel()
+       {
+               task_pool.stop();
+
+               /* Release kernels */
+               program_data_init.release();
+               program_scene_intersect.release();
+               program_lamp_emission.release();
+               program_queue_enqueue.release();
+               program_background_buffer_update.release();
+               program_shader_eval.release();
+               program_holdout_emission_blurring_pathtermination_ao.release();
+               program_direct_lighting.release();
+               program_shadow_blocked.release();
+               program_next_iteration_setup.release();
+               program_sum_all_radiance.release();
+
+               /* Release global memory */
+               release_mem_object_safe(rng_coop);
+               release_mem_object_safe(throughput_coop);
+               release_mem_object_safe(L_transparent_coop);
+               release_mem_object_safe(PathRadiance_coop);
+               release_mem_object_safe(Ray_coop);
+               release_mem_object_safe(PathState_coop);
+               release_mem_object_safe(Intersection_coop);
+               release_mem_object_safe(kgbuffer);
+               release_mem_object_safe(sd);
+               release_mem_object_safe(sd_DL_shadow);
+               release_mem_object_safe(ray_state);
+               release_mem_object_safe(AOAlpha_coop);
+               release_mem_object_safe(AOBSDF_coop);
+               release_mem_object_safe(AOLightRay_coop);
+               release_mem_object_safe(BSDFEval_coop);
+               release_mem_object_safe(ISLamp_coop);
+               release_mem_object_safe(LightRay_coop);
+               release_mem_object_safe(Intersection_coop_shadow);
+#ifdef WITH_CYCLES_DEBUG
+               release_mem_object_safe(debugdata_coop);
+#endif
+               release_mem_object_safe(use_queues_flag);
+               release_mem_object_safe(Queue_data);
+               release_mem_object_safe(Queue_index);
+               release_mem_object_safe(work_array);
+#ifdef __WORK_STEALING__
+               release_mem_object_safe(work_pool_wgs);
+#endif
+               release_mem_object_safe(per_sample_output_buffers);
+
+               if(hostRayStateArray != NULL) {
+                       free(hostRayStateArray);
+               }
+       }
+
+       void path_trace(DeviceTask *task,
+                       SplitRenderTile& rtile,
+                       int2 max_render_feasible_tile_size)
+       {
+               /* cast arguments to cl types */
+               cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+               cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
+               cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
+               cl_int d_x = rtile.x;
+               cl_int d_y = rtile.y;
+               cl_int d_w = rtile.w;
+               cl_int d_h = rtile.h;
+               cl_int d_offset = rtile.offset;
+               cl_int d_stride = rtile.stride;
+
+               /* Make sure that set render feasible tile size is a multiple of local
+                * work size dimensions.
+                */
+               assert(max_render_feasible_tile_size.x % SPLIT_KERNEL_LOCAL_SIZE_X == 0);
+               assert(max_render_feasible_tile_size.y % SPLIT_KERNEL_LOCAL_SIZE_Y == 0);
+
+               size_t global_size[2];
+               size_t local_size[2] = {SPLIT_KERNEL_LOCAL_SIZE_X,
+                                       SPLIT_KERNEL_LOCAL_SIZE_Y};
+
+               /* Set the range of samples to be processed for every ray in
+                * path-regeneration logic.
+                */
+               cl_int start_sample = rtile.start_sample;
+               cl_int end_sample = rtile.start_sample + rtile.num_samples;
+               cl_int num_samples = rtile.num_samples;
+
+#ifdef __WORK_STEALING__
+               global_size[0] = (((d_w - 1) / local_size[0]) + 1) * local_size[0];
+               global_size[1] = (((d_h - 1) / local_size[1]) + 1) * local_size[1];
+               unsigned int num_parallel_samples = 1;
+#else
+               global_size[1] = (((d_h - 1) / local_size[1]) + 1) * local_size[1];
+               unsigned int num_threads = max_render_feasible_tile_size.x *
+                                          max_render_feasible_tile_size.y;
+               unsigned int num_tile_columns_possible = num_threads / global_size[1];
+               /* Estimate number of parallel samples that can be
+                * processed in parallel.
+                */
+               unsigned int num_parallel_samples = min(num_tile_columns_possible / d_w,
+                                                       rtile.num_samples);
+               /* Wavefront size in AMD is 64.
+                * TODO(sergey): What about other platforms?
+                */
+               if(num_parallel_samples >= 64) {
+                       /* TODO(sergey): Could use generic round-up here. */
+                       num_parallel_samples = (num_parallel_samples / 64) * 64;
+               }
+               assert(num_parallel_samples != 0);
+
+               global_size[0] = d_w * num_parallel_samples;
+#endif  /* __WORK_STEALING__ */
+
+               assert(global_size[0] * global_size[1] <=
+                      max_render_feasible_tile_size.x * max_render_feasible_tile_size.y);
+
+               /* Allocate all required global memory once. */
+               if(first_tile) {
+                       size_t num_global_elements = max_render_feasible_tile_size.x *
+                                                    max_render_feasible_tile_size.y;
+                       /* TODO(sergey): This will actually over-allocate if
+                        * particular kernel does not support multiclosure.
+                        */
+                       size_t shaderdata_size = get_shader_data_size(current_max_closure);
+
+#ifdef __WORK_STEALING__
+                       /* Calculate max groups */
+                       size_t max_global_size[2];
+                       size_t tile_x = max_render_feasible_tile_size.x;
+                       size_t tile_y = max_render_feasible_tile_size.y;
+                       max_global_size[0] = (((tile_x - 1) / local_size[0]) + 1) * local_size[0];
+                       max_global_size[1] = (((tile_y - 1) / local_size[1]) + 1) * local_size[1];
+                       max_work_groups = (max_global_size[0] * max_global_size[1]) /
+                                         (local_size[0] * local_size[1]);
+                       /* Allocate work_pool_wgs memory. */
+                       work_pool_wgs = mem_alloc(max_work_groups * sizeof(unsigned int));
+#endif  /* __WORK_STEALING__ */
+
+                       /* Allocate queue_index memory only once. */
+                       Queue_index = mem_alloc(NUM_QUEUES * sizeof(int));
+                       use_queues_flag = mem_alloc(sizeof(char));
+                       kgbuffer = mem_alloc(get_KernelGlobals_size());
+
+                       /* Create global buffers for ShaderData. */
+                       sd = mem_alloc(num_global_elements * shaderdata_size);
+                       sd_DL_shadow = mem_alloc(num_global_elements * 2 * shaderdata_size);
+
+                       /* Creation of global memory buffers which are shared among
+                        * the kernels.
+                        */
+                       rng_coop = mem_alloc(num_global_elements * sizeof(RNG));
+                       throughput_coop = mem_alloc(num_global_elements * sizeof(float3));
+                       L_transparent_coop = mem_alloc(num_global_elements * sizeof(float));
+                       PathRadiance_coop = mem_alloc(num_global_elements * sizeof(PathRadiance));
+                       Ray_coop = mem_alloc(num_global_elements * sizeof(Ray));
+                       PathState_coop = mem_alloc(num_global_elements * sizeof(PathState));
+                       Intersection_coop = mem_alloc(num_global_elements * sizeof(Intersection));
+                       AOAlpha_coop = mem_alloc(num_global_elements * sizeof(float3));
+                       AOBSDF_coop = mem_alloc(num_global_elements * sizeof(float3));
+                       AOLightRay_coop = mem_alloc(num_global_elements * sizeof(Ray));
+                       BSDFEval_coop = mem_alloc(num_global_elements * sizeof(BsdfEval));
+                       ISLamp_coop = mem_alloc(num_global_elements * sizeof(int));
+                       LightRay_coop = mem_alloc(num_global_elements * sizeof(Ray));
+                       Intersection_coop_shadow = mem_alloc(2 * num_global_elements * sizeof(Intersection));
+
+#ifdef WITH_CYCLES_DEBUG
+                       debugdata_coop = mem_alloc(num_global_elements * sizeof(DebugData));
+#endif
+
+                       ray_state = mem_alloc(num_global_elements * sizeof(char));
+
+                       hostRayStateArray = (char *)calloc(num_global_elements, sizeof(char));
+                       assert(hostRayStateArray != NULL && "Can't create hostRayStateArray memory");
+
+                       Queue_data = mem_alloc(num_global_elements * (NUM_QUEUES * sizeof(int)+sizeof(int)));
+                       work_array = mem_alloc(num_global_elements * sizeof(unsigned int));
+                       per_sample_output_buffers = mem_alloc(num_global_elements *
+                                                             per_thread_output_buffer_size);
+               }
+
+               cl_int dQueue_size = global_size[0] * global_size[1];
+
+               cl_uint start_arg_index =
+                       kernel_set_args(program_data_init(),
+                                       0,
+                                       kgbuffer,
+                                       sd_DL_shadow,
+                                       d_data,
+                                       per_sample_output_buffers,
+                                       d_rng_state,
+                                       rng_coop,
+                                       throughput_coop,
+                                       L_transparent_coop,
+                                       PathRadiance_coop,
+                                       Ray_coop,
+                                       PathState_coop,
+                                       Intersection_coop_shadow,
+                                       ray_state);
+
+/* TODO(sergey): Avoid map lookup here. */
+#define KERNEL_TEX(type, ttype, name) \
+       set_kernel_arg_mem(program_data_init(), &start_arg_index, #name);
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+
+               start_arg_index +=
+                       kernel_set_args(program_data_init(),
+                                       start_arg_index,
+                                       start_sample,
+                                       d_x,
+                                       d_y,
+                                       d_w,
+                                       d_h,
+                                       d_offset,
+                                       d_stride,
+                                       rtile.rng_state_offset_x,
+                                       rtile.rng_state_offset_y,
+                                       rtile.buffer_rng_state_stride,
+                                       Queue_data,
+                                       Queue_index,
+                                       dQueue_size,
+                                       use_queues_flag,
+                                       work_array,
+#ifdef __WORK_STEALING__
+                                       work_pool_wgs,
+                                       num_samples,
+#endif
+#ifdef WITH_CYCLES_DEBUG
+                                       debugdata_coop,
+#endif
+                                       num_parallel_samples);
+
+               kernel_set_args(program_scene_intersect(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               rng_coop,
+                               Ray_coop,
+                               PathState_coop,
+                               Intersection_coop,
+                               ray_state,
+                               d_w,
+                               d_h,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size,
+                               use_queues_flag,
+#ifdef WITH_CYCLES_DEBUG
+                               debugdata_coop,
+#endif
+                               num_parallel_samples);
+
+               kernel_set_args(program_lamp_emission(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               throughput_coop,
+                               PathRadiance_coop,
+                               Ray_coop,
+                               PathState_coop,
+                               Intersection_coop,
+                               ray_state,
+                               d_w,
+                               d_h,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size,
+                               use_queues_flag,
+                               num_parallel_samples);
+
+               kernel_set_args(program_queue_enqueue(),
+                               0,
+                               Queue_data,
+                               Queue_index,
+                               ray_state,
+                               dQueue_size);
+
+               kernel_set_args(program_background_buffer_update(),
+                                0,
+                                kgbuffer,
+                                d_data,
+                                per_sample_output_buffers,
+                                d_rng_state,
+                                rng_coop,
+                                throughput_coop,
+                                PathRadiance_coop,
+                                Ray_coop,
+                                PathState_coop,
+                                L_transparent_coop,
+                                ray_state,
+                                d_w,
+                                d_h,
+                                d_x,
+                                d_y,
+                                d_stride,
+                                rtile.rng_state_offset_x,
+                                rtile.rng_state_offset_y,
+                                rtile.buffer_rng_state_stride,
+                                work_array,
+                                Queue_data,
+                                Queue_index,
+                                dQueue_size,
+                                end_sample,
+                                start_sample,
+#ifdef __WORK_STEALING__
+                                work_pool_wgs,
+                                num_samples,
+#endif
+#ifdef WITH_CYCLES_DEBUG
+                                debugdata_coop,
+#endif
+                                num_parallel_samples);
+
+               kernel_set_args(program_shader_eval(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               sd,
+                               rng_coop,
+                               Ray_coop,
+                               PathState_coop,
+                               Intersection_coop,
+                               ray_state,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size);
+
+               kernel_set_args(program_holdout_emission_blurring_pathtermination_ao(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               sd,
+                               per_sample_output_buffers,
+                               rng_coop,
+                               throughput_coop,
+                               L_transparent_coop,
+                               PathRadiance_coop,
+                               PathState_coop,
+                               Intersection_coop,
+                               AOAlpha_coop,
+                               AOBSDF_coop,
+                               AOLightRay_coop,
+                               d_w,
+                               d_h,
+                               d_x,
+                               d_y,
+                               d_stride,
+                               ray_state,
+                               work_array,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size,
+#ifdef __WORK_STEALING__
+                               start_sample,
+#endif
+                               num_parallel_samples);
+
+               kernel_set_args(program_direct_lighting(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               sd,
+                               rng_coop,
+                               PathState_coop,
+                               ISLamp_coop,
+                               LightRay_coop,
+                               BSDFEval_coop,
+                               ray_state,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size);
+
+               kernel_set_args(program_shadow_blocked(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               PathState_coop,
+                               LightRay_coop,
+                               AOLightRay_coop,
+                               ray_state,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size);
+
+               kernel_set_args(program_next_iteration_setup(),
+                               0,
+                               kgbuffer,
+                               d_data,
+                               sd,
+                               rng_coop,
+                               throughput_coop,
+                               PathRadiance_coop,
+                               Ray_coop,
+                               PathState_coop,
+                               LightRay_coop,
+                               ISLamp_coop,
+                               BSDFEval_coop,
+                               AOLightRay_coop,
+                               AOBSDF_coop,
+                               AOAlpha_coop,
+                               ray_state,
+                               Queue_data,
+                               Queue_index,
+                               dQueue_size,
+                               use_queues_flag);
+
+               kernel_set_args(program_sum_all_radiance(),
+                               0,
+                               d_data,
+                               d_buffer,
+                               per_sample_output_buffers,
+                               num_parallel_samples,
+                               d_w,
+                               d_h,
+                               d_stride,
+                               rtile.buffer_offset_x,
+                               rtile.buffer_offset_y,
+                               rtile.buffer_rng_state_stride,
+                               start_sample);
+
+               /* Macro for Enqueuing split kernels. */
+#define GLUE(a, b) a ## b
+#define ENQUEUE_SPLIT_KERNEL(kernelName, globalSize, localSize) \
+               { \
+                       ciErr = clEnqueueNDRangeKernel(cqCommandQueue, \
+                                                      GLUE(program_, \
+                                                           kernelName)(), \
+                                                      2, \
+                                                      NULL, \
+                                                      globalSize, \
+                                                      localSize, \
+                                                      0, \
+                                                      NULL, \
+                                                      NULL); \
+                       opencl_assert_err(ciErr, "clEnqueueNDRangeKernel"); \
+                       if(ciErr != CL_SUCCESS) { \
+                               string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()", \
+                                                              clewErrorString(ciErr)); \
+                               opencl_error(message); \
+                               return; \
+                       } \
+               } (void) 0
+
+               /* Enqueue ckPathTraceKernel_data_init kernel. */
+               ENQUEUE_SPLIT_KERNEL(data_init, global_size, local_size);
+               bool activeRaysAvailable = true;
+
+               /* Record number of time host intervention has been made */
+               unsigned int numHostIntervention = 0;
+               unsigned int numNextPathIterTimes = PathIteration_times;
+               bool canceled = false;
+               while(activeRaysAvailable) {
+                       /* Twice the global work size of other kernels for
+                        * ckPathTraceKernel_shadow_blocked_direct_lighting. */
+                       size_t global_size_shadow_blocked[2];
+                       global_size_shadow_blocked[0] = global_size[0] * 2;
+                       global_size_shadow_blocked[1] = global_size[1];
+
+                       /* Do path-iteration in host [Enqueue Path-iteration kernels. */
+                       for(int PathIter = 0; PathIter < PathIteration_times; PathIter++) {
+                               ENQUEUE_SPLIT_KERNEL(scene_intersect, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(lamp_emission, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(background_buffer_update, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(shader_eval, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(holdout_emission_blurring_pathtermination_ao, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
+                               ENQUEUE_SPLIT_KERNEL(shadow_blocked, global_size_shadow_blocked, local_size);
+                               ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size);
+
+                               if(task->get_cancel()) {
+                                       canceled = true;
+                                       break;
+                               }
+                       }
+
+                       /* Read ray-state into Host memory to decide if we should exit
+                        * path-iteration in host.
+                        */
+                       ciErr = clEnqueueReadBuffer(cqCommandQueue,
+                                                   ray_state,
+                                                   CL_TRUE,
+                                                   0,
+                                                   global_size[0] * global_size[1] * sizeof(char),
+                                                   hostRayStateArray,
+                                                   0,
+                                                   NULL,
+                                                   NULL);
+                       assert(ciErr == CL_SUCCESS);
+
+                       activeRaysAvailable = false;
+
+                       for(int rayStateIter = 0;
+                           rayStateIter < global_size[0] * global_size[1];
+                           ++rayStateIter)
+                       {
+                               if(int8_t(hostRayStateArray[rayStateIter]) != RAY_INACTIVE) {
+                                       /* Not all rays are RAY_INACTIVE. */
+                                       activeRaysAvailable = true;
+                                       break;
+                               }
+                       }
+
+                       if(activeRaysAvailable) {
+                               numHostIntervention++;
+                               PathIteration_times = PATH_ITER_INC_FACTOR;
+                               /* Host intervention done before all rays become RAY_INACTIVE;
+                                * Set do more initial iterations for the next tile.
+                                */
+                               numNextPathIterTimes += PATH_ITER_INC_FACTOR;
+                       }
+
+                       if(task->get_cancel()) {
+                               canceled = true;
+                               break;
+                       }
+               }
+
+               /* Execute SumALLRadiance kernel to accumulate radiance calculated in
+                * per_sample_output_buffers into RenderTile's output buffer.
+                */
+               if(!canceled) {
+                       size_t sum_all_radiance_local_size[2] = {16, 16};
+                       size_t sum_all_radiance_global_size[2];
+                       sum_all_radiance_global_size[0] =
+                               (((d_w - 1) / sum_all_radiance_local_size[0]) + 1) *
+                               sum_all_radiance_local_size[0];
+                       sum_all_radiance_global_size[1] =
+                               (((d_h - 1) / sum_all_radiance_local_size[1]) + 1) *
+                               sum_all_radiance_local_size[1];
+                       ENQUEUE_SPLIT_KERNEL(sum_all_radiance,
+                                            sum_all_radiance_global_size,
+                                            sum_all_radiance_local_size);
+               }
+
+#undef ENQUEUE_SPLIT_KERNEL
+#undef GLUE
+
+               if(numHostIntervention == 0) {
+                       /* This means that we are executing kernel more than required
+                        * Must avoid this for the next sample/tile.
+                        */
+                       PathIteration_times = ((numNextPathIterTimes - PATH_ITER_INC_FACTOR) <= 0) ?
+                       PATH_ITER_INC_FACTOR : numNextPathIterTimes - PATH_ITER_INC_FACTOR;
+               }
+               else {
+                       /* Number of path-iterations done for this tile is set as
+                        * Initial path-iteration times for the next tile
+                        */
+                       PathIteration_times = numNextPathIterTimes;
+               }
+
+               first_tile = false;
+       }
+
+       /* Calculates the amount of memory that has to be always
+        * allocated in order for the split kernel to function.
+        * This memory is tile/scene-property invariant (meaning,
+        * the value returned by this function does not depend
+        * on the user set tile size or scene properties.
+        */
+       size_t get_invariable_mem_allocated()
+       {
+               size_t total_invariable_mem_allocated = 0;
+               size_t KernelGlobals_size = 0;
+
+               KernelGlobals_size = get_KernelGlobals_size();
+
+               total_invariable_mem_allocated += KernelGlobals_size; /* KernelGlobals size */
+               total_invariable_mem_allocated += NUM_QUEUES * sizeof(unsigned int); /* Queue index size */
+               total_invariable_mem_allocated += sizeof(char); /* use_queues_flag size */
+
+               return total_invariable_mem_allocated;
+       }
+
+       /* Calculate the memory that has-to-be/has-been allocated for
+        * the split kernel to function.
+        */
+       size_t get_tile_specific_mem_allocated(const int2 tile_size)
+       {
+               size_t tile_specific_mem_allocated = 0;
+
+               /* Get required tile info */
+               unsigned int user_set_tile_w = tile_size.x;
+               unsigned int user_set_tile_h = tile_size.y;
+
+#ifdef __WORK_STEALING__
+               /* Calculate memory to be allocated for work_pools in
+                * case of work_stealing.
+                */
+               size_t max_global_size[2];
+               size_t max_num_work_pools = 0;
+               max_global_size[0] =
+                       (((user_set_tile_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_X;
+               max_global_size[1] =
+                       (((user_set_tile_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+               max_num_work_pools =
+                       (max_global_size[0] * max_global_size[1]) /
+                       (SPLIT_KERNEL_LOCAL_SIZE_X * SPLIT_KERNEL_LOCAL_SIZE_Y);
+               tile_specific_mem_allocated += max_num_work_pools * sizeof(unsigned int);
+#endif
+
+               tile_specific_mem_allocated +=
+                       user_set_tile_w * user_set_tile_h * per_thread_output_buffer_size;
+               tile_specific_mem_allocated +=
+                       user_set_tile_w * user_set_tile_h * sizeof(RNG);
+
+               return tile_specific_mem_allocated;
+       }
+
+       /* Calculates the texture memories and KernelData (d_data) memory
+        * that has been allocated.
+        */
+       size_t get_scene_specific_mem_allocated(cl_mem d_data)
+       {
+               size_t scene_specific_mem_allocated = 0;
+               /* Calculate texture memories. */
+#define KERNEL_TEX(type, ttype, name) \
+       scene_specific_mem_allocated += get_tex_size(#name);
+#include "kernel_textures.h"
+#undef KERNEL_TEX
+               size_t d_data_size;
+               ciErr = clGetMemObjectInfo(d_data,
+                                          CL_MEM_SIZE,
+                                          sizeof(d_data_size),
+                                          &d_data_size,
+                                          NULL);
+               assert(ciErr == CL_SUCCESS && "Can't get d_data mem object info");
+               scene_specific_mem_allocated += d_data_size;
+               return scene_specific_mem_allocated;
+       }
+
+       /* Calculate the memory required for one thread in split kernel. */
+       size_t get_per_thread_memory()
+       {
+               size_t shaderdata_size = 0;
+               /* TODO(sergey): This will actually over-allocate if
+                * particular kernel does not support multiclosure.
+                */
+               shaderdata_size = get_shader_data_size(current_max_closure);
+               size_t retval = sizeof(RNG)
+                       + sizeof(float3)          /* Throughput size */
+                       + sizeof(float)           /* L transparent size */
+                       + sizeof(char)            /* Ray state size */
+                       + sizeof(unsigned int)    /* Work element size */
+                       + sizeof(int)             /* ISLamp_size */
+                       + sizeof(PathRadiance) + sizeof(Ray) + sizeof(PathState)
+                       + sizeof(Intersection)    /* Overall isect */
+                       + sizeof(Intersection)    /* Instersection_coop_AO */
+                       + sizeof(Intersection)    /* Intersection coop DL */
+                       + shaderdata_size         /* Overall ShaderData */
+                       + (shaderdata_size * 2)   /* ShaderData : DL and shadow */
+                       + sizeof(Ray) + sizeof(BsdfEval)
+                       + sizeof(float3)          /* AOAlpha size */
+                       + sizeof(float3)          /* AOBSDF size */
+                       + sizeof(Ray)
+                       + (sizeof(int) * NUM_QUEUES)
+                       + per_thread_output_buffer_size;
+               return retval;
+       }
+
+       /* Considers the total memory available in the device and
+        * and returns the maximum global work size possible.
+        */
+       size_t get_feasible_global_work_size(int2 tile_size, cl_mem d_data)
+       {
+               /* Calculate invariably allocated memory. */
+               size_t invariable_mem_allocated = get_invariable_mem_allocated();
+               /* Calculate tile specific allocated memory. */
+               size_t tile_specific_mem_allocated =
+                       get_tile_specific_mem_allocated(tile_size);
+               /* Calculate scene specific allocated memory. */
+               size_t scene_specific_mem_allocated =
+                       get_scene_specific_mem_allocated(d_data);
+               /* Calculate total memory available for the threads in global work size. */
+               size_t available_memory = total_allocatable_memory
+                       - invariable_mem_allocated
+                       - tile_specific_mem_allocated
+                       - scene_specific_mem_allocated
+                       - DATA_ALLOCATION_MEM_FACTOR;
+               size_t per_thread_memory_required = get_per_thread_memory();
+               return (available_memory / per_thread_memory_required);
+       }
+
+       /* Checks if the device has enough memory to render the whole tile;
+        * If not, we should split single tile into multiple tiles of small size
+        * and process them all.
+        */
+       bool need_to_split_tile(unsigned int d_w,
+                               unsigned int d_h,
+                               int2 max_render_feasible_tile_size)
+       {
+               size_t global_size_estimate[2];
+               /* TODO(sergey): Such round-ups are in quite few places, need to replace
+                * them with an utility macro.
+                */
+               global_size_estimate[0] =
+                       (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_X;
+               global_size_estimate[1] =
+                       (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+               if((global_size_estimate[0] * global_size_estimate[1]) >
+                  (max_render_feasible_tile_size.x * max_render_feasible_tile_size.y))
+               {
+                       return true;
+               }
+               else {
+                       return false;
+               }
+       }
+
+       /* Considers the scene properties, global memory available in the device
+        * and returns a rectanglular tile dimension (approx the maximum)
+        * that should render on split kernel.
+        */
+       int2 get_max_render_feasible_tile_size(size_t feasible_global_work_size)
+       {
+               int2 max_render_feasible_tile_size;
+               int square_root_val = (int)sqrt(feasible_global_work_size);
+               max_render_feasible_tile_size.x = square_root_val;
+               max_render_feasible_tile_size.y = square_root_val;
+               /* Ciel round-off max_render_feasible_tile_size. */
+               int2 ceil_render_feasible_tile_size;
+               ceil_render_feasible_tile_size.x =
+                       (((max_render_feasible_tile_size.x - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_X;
+               ceil_render_feasible_tile_size.y =
+                       (((max_render_feasible_tile_size.y - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+               if(ceil_render_feasible_tile_size.x * ceil_render_feasible_tile_size.y <=
+                  feasible_global_work_size)
+               {
+                       return ceil_render_feasible_tile_size;
+               }
+               /* Floor round-off max_render_feasible_tile_size. */
+               int2 floor_render_feasible_tile_size;
+               floor_render_feasible_tile_size.x =
+                       (max_render_feasible_tile_size.x / SPLIT_KERNEL_LOCAL_SIZE_X) *
+                       SPLIT_KERNEL_LOCAL_SIZE_X;
+               floor_render_feasible_tile_size.y =
+                       (max_render_feasible_tile_size.y / SPLIT_KERNEL_LOCAL_SIZE_Y) *
+                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+               return floor_render_feasible_tile_size;
+       }
+
+       /* Try splitting the current tile into multiple smaller
+        * almost-square-tiles.
+        */
+       int2 get_split_tile_size(RenderTile rtile,
+                                int2 max_render_feasible_tile_size)
+       {
+               int2 split_tile_size;
+               int num_global_threads = max_render_feasible_tile_size.x *
+                                        max_render_feasible_tile_size.y;
+               int d_w = rtile.w;
+               int d_h = rtile.h;
+               /* Ceil round off d_w and d_h */
+               d_w = (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_X;
+               d_h = (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+               while(d_w * d_h > num_global_threads) {
+                       /* Halve the longer dimension. */
+                       if(d_w >= d_h) {
+                               d_w = d_w / 2;
+                               d_w = (((d_w - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                                       SPLIT_KERNEL_LOCAL_SIZE_X;
+                       }
+                       else {
+                               d_h = d_h / 2;
+                               d_h = (((d_h - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                                       SPLIT_KERNEL_LOCAL_SIZE_Y;
+                       }
+               }
+               split_tile_size.x = d_w;
+               split_tile_size.y = d_h;
+               return split_tile_size;
+       }
+
+       /* Splits existing tile into multiple tiles of tile size split_tile_size. */
+       vector<SplitRenderTile> split_tiles(RenderTile rtile, int2 split_tile_size)
+       {
+               vector<SplitRenderTile> to_path_trace_rtile;
+               int d_w = rtile.w;
+               int d_h = rtile.h;
+               int num_tiles_x = (((d_w - 1) / split_tile_size.x) + 1);
+               int num_tiles_y = (((d_h - 1) / split_tile_size.y) + 1);
+               /* Buffer and rng_state offset calc. */
+               size_t offset_index = rtile.offset + (rtile.x + rtile.y * rtile.stride);
+               size_t offset_x = offset_index % rtile.stride;
+               size_t offset_y = offset_index / rtile.stride;
+               /* Resize to_path_trace_rtile. */
+               to_path_trace_rtile.resize(num_tiles_x * num_tiles_y);
+               for(int tile_iter_y = 0; tile_iter_y < num_tiles_y; tile_iter_y++) {
+                       for(int tile_iter_x = 0; tile_iter_x < num_tiles_x; tile_iter_x++) {
+                               int rtile_index = tile_iter_y * num_tiles_x + tile_iter_x;
+                               to_path_trace_rtile[rtile_index].rng_state_offset_x = offset_x + tile_iter_x * split_tile_size.x;
+                               to_path_trace_rtile[rtile_index].rng_state_offset_y = offset_y + tile_iter_y * split_tile_size.y;
+                               to_path_trace_rtile[rtile_index].buffer_offset_x = offset_x + tile_iter_x * split_tile_size.x;
+                               to_path_trace_rtile[rtile_index].buffer_offset_y = offset_y + tile_iter_y * split_tile_size.y;
+                               to_path_trace_rtile[rtile_index].start_sample = rtile.start_sample;
+                               to_path_trace_rtile[rtile_index].num_samples = rtile.num_samples;
+                               to_path_trace_rtile[rtile_index].sample = rtile.sample;
+                               to_path_trace_rtile[rtile_index].resolution = rtile.resolution;
+                               to_path_trace_rtile[rtile_index].offset = rtile.offset;
+                               to_path_trace_rtile[rtile_index].buffers = rtile.buffers;
+                               to_path_trace_rtile[rtile_index].buffer = rtile.buffer;
+                               to_path_trace_rtile[rtile_index].rng_state = rtile.rng_state;
+                               to_path_trace_rtile[rtile_index].x = rtile.x + (tile_iter_x * split_tile_size.x);
+                               to_path_trace_rtile[rtile_index].y = rtile.y + (tile_iter_y * split_tile_size.y);
+                               to_path_trace_rtile[rtile_index].buffer_rng_state_stride = rtile.stride;
+                               /* Fill width and height of the new render tile. */
+                               to_path_trace_rtile[rtile_index].w = (tile_iter_x == (num_tiles_x - 1)) ?
+                                       (d_w - (tile_iter_x * split_tile_size.x)) /* Border tile */
+                                       : split_tile_size.x;
+                               to_path_trace_rtile[rtile_index].h = (tile_iter_y == (num_tiles_y - 1)) ?
+                                       (d_h - (tile_iter_y * split_tile_size.y)) /* Border tile */
+                                       : split_tile_size.y;
+                               to_path_trace_rtile[rtile_index].stride = to_path_trace_rtile[rtile_index].w;
+                       }
+               }
+               return to_path_trace_rtile;
+       }
+
+       void thread_run(DeviceTask *task)
+       {
+               if(task->type == DeviceTask::FILM_CONVERT) {
+                       film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
+               }
+               else if(task->type == DeviceTask::SHADER) {
+                       shader(*task);
+               }
+               else if(task->type == DeviceTask::PATH_TRACE) {
+                       RenderTile tile;
+                       bool initialize_data_and_check_render_feasibility = false;
+                       bool need_to_split_tiles_further = false;
+                       int2 max_render_feasible_tile_size;
+                       size_t feasible_global_work_size;
+                       const int2 tile_size = task->requested_tile_size;
+                       /* Keep rendering tiles until done. */
+                       while(task->acquire_tile(this, tile)) {
+                               if(!initialize_data_and_check_render_feasibility) {
+                                       /* Initialize data. */
+                                       /* Calculate per_thread_output_buffer_size. */
+                                       size_t output_buffer_size = 0;
+                                       ciErr = clGetMemObjectInfo((cl_mem)tile.buffer,
+                                                                  CL_MEM_SIZE,
+                                                                  sizeof(output_buffer_size),
+                                                                  &output_buffer_size,
+                                                                  NULL);
+                                       assert(ciErr == CL_SUCCESS && "Can't get tile.buffer mem object info");
+                                       /* This value is different when running on AMD and NV. */
+                                       if(background) {
+                                               /* In offline render the number of buffer elements
+                                                * associated with tile.buffer is the current tile size.
+                                                */
+                                               per_thread_output_buffer_size =
+                                                       output_buffer_size / (tile.w * tile.h);
+                                       }
+                                       else {
+                                               /* interactive rendering, unlike offline render, the number of buffer elements
+                                                * associated with tile.buffer is the entire viewport size.
+                                                */
+                                               per_thread_output_buffer_size =
+                                                       output_buffer_size / (tile.buffers->params.width *
+                                                                             tile.buffers->params.height);
+                                       }
+                                       /* Check render feasibility. */
+                                       feasible_global_work_size = get_feasible_global_work_size(
+                                               tile_size,
+                                               CL_MEM_PTR(const_mem_map["__data"]->device_pointer));
+                                       max_render_feasible_tile_size =
+                                               get_max_render_feasible_tile_size(
+                                                       feasible_global_work_size);
+                                       need_to_split_tiles_further =
+                                               need_to_split_tile(tile_size.x,
+                                                                  tile_size.y,
+                                                                  max_render_feasible_tile_size);
+                                       initialize_data_and_check_render_feasibility = true;
+                               }
+                               if(need_to_split_tiles_further) {
+                                       int2 split_tile_size =
+                                               get_split_tile_size(tile,
+                                                                   max_render_feasible_tile_size);
+                                       vector<SplitRenderTile> to_path_trace_render_tiles =
+                                               split_tiles(tile, split_tile_size);
+                                       /* Print message to console */
+                                       if(background && (to_path_trace_render_tiles.size() > 1)) {
+                                               fprintf(stderr, "Message : Tiles need to be split "
+                                                       "further inside path trace (due to insufficient "
+                                                       "device-global-memory for split kernel to "
+                                                       "function) \n"
+                                                       "The current tile of dimensions %dx%d is split "
+                                                       "into tiles of dimension %dx%d for render \n",
+                                                       tile.w, tile.h,
+                                                       split_tile_size.x,
+                                                       split_tile_size.y);
+                                       }
+                                       /* Process all split tiles. */
+                                       for(int tile_iter = 0;
+                                           tile_iter < to_path_trace_render_tiles.size();
+                                           ++tile_iter)
+                                       {
+                                               path_trace(task,
+                                                          to_path_trace_render_tiles[tile_iter],
+                                                          max_render_feasible_tile_size);
+                                       }
+                               }
+                               else {
+                                       /* No splitting required; process the entire tile at once. */
+                                       /* Render feasible tile size is user-set-tile-size itself. */
+                                       max_render_feasible_tile_size.x =
+                                               (((tile_size.x - 1) / SPLIT_KERNEL_LOCAL_SIZE_X) + 1) *
+                                               SPLIT_KERNEL_LOCAL_SIZE_X;
+                                       max_render_feasible_tile_size.y =
+                                               (((tile_size.y - 1) / SPLIT_KERNEL_LOCAL_SIZE_Y) + 1) *
+                                               SPLIT_KERNEL_LOCAL_SIZE_Y;
+                                       /* buffer_rng_state_stride is stride itself. */
+                                       SplitRenderTile split_tile(tile);
+                                       split_tile.buffer_rng_state_stride = tile.stride;
+                                       path_trace(task, split_tile, max_render_feasible_tile_size);
+                               }
+                               tile.sample = tile.start_sample + tile.num_samples;
+
+                               /* Complete kernel execution before release tile. */
+                               /* This helps in multi-device render;
+                                * The device that reaches the critical-section function
+                                * release_tile waits (stalling other devices from entering
+                                * release_tile) for all kernels to complete. If device1 (a
+                                * slow-render device) reaches release_tile first then it would
+                                * stall device2 (a fast-render device) from proceeding to render
+                                * next tile.
+                                */
+                               clFinish(cqCommandQueue);
+
+                               task->release_tile(tile);
+                       }
+               }
+       }
+
+protected:
+       cl_mem mem_alloc(size_t bufsize, cl_mem_flags mem_flag = CL_MEM_READ_WRITE)
+       {
+               cl_mem ptr;
+               assert(bufsize != 0);
+               ptr = clCreateBuffer(cxContext, mem_flag, bufsize, NULL, &ciErr);
+               opencl_assert_err(ciErr, "clCreateBuffer");
+               return ptr;
+       }
+
+       /* ** Those guys are for workign around some compiler-specific bugs ** */
+
+       string build_options_for_base_program(
+               const DeviceRequestedFeatures& requested_features)
+       {
+               return requested_features.get_build_options();
+       }
+};
+
+Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, bool background)
+{
+       return new OpenCLDeviceSplitKernel(info, stats, background);
+}
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_OPENCL */
diff --git a/intern/cycles/device/opencl/opencl_util.cpp b/intern/cycles/device/opencl/opencl_util.cpp
new file mode 100644 (file)
index 0000000..9650455
--- /dev/null
@@ -0,0 +1,761 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifdef WITH_OPENCL
+
+#include "opencl.h"
+
+#include "util_logging.h"
+#include "util_path.h"
+#include "util_time.h"
+
+using std::cerr;
+using std::endl;
+
+CCL_NAMESPACE_BEGIN
+
+OpenCLCache::Slot::ProgramEntry::ProgramEntry()
+ : program(NULL),
+   mutex(NULL)
+{
+}
+
+OpenCLCache::Slot::ProgramEntry::ProgramEntry(const ProgramEntry& rhs)
+ : program(rhs.program),
+   mutex(NULL)
+{
+}
+
+OpenCLCache::Slot::ProgramEntry::~ProgramEntry()
+{
+       delete mutex;
+}
+
+OpenCLCache::Slot::Slot()
+ : context_mutex(NULL),
+   context(NULL)
+{
+}
+
+OpenCLCache::Slot::Slot(const Slot& rhs)
+ : context_mutex(NULL),
+   context(NULL),
+   programs(rhs.programs)
+{
+}
+
+OpenCLCache::Slot::~Slot()
+{
+       delete context_mutex;
+}
+
+OpenCLCache& OpenCLCache::global_instance()
+{
+       static OpenCLCache instance;
+       return instance;
+}
+
+cl_context OpenCLCache::get_context(cl_platform_id platform,
+                                    cl_device_id device,
+                                    thread_scoped_lock& slot_locker)
+{
+       assert(platform != NULL);
+
+       OpenCLCache& self = global_instance();
+
+       thread_scoped_lock cache_lock(self.cache_lock);
+
+       pair<CacheMap::iterator,bool> ins = self.cache.insert(
+               CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
+
+       Slot &slot = ins.first->second;
+
+       /* create slot lock only while holding cache lock */
+       if(!slot.context_mutex)
+               slot.context_mutex = new thread_mutex;
+
+       /* need to unlock cache before locking slot, to allow store to complete */
+       cache_lock.unlock();
+
+       /* lock the slot */
+       slot_locker = thread_scoped_lock(*slot.context_mutex);
+
+       /* If the thing isn't cached */
+       if(slot.context == NULL) {
+               /* return with the caller's lock holder holding the slot lock */
+               return NULL;
+       }
+
+       /* the item was already cached, release the slot lock */
+       slot_locker.unlock();
+
+       cl_int ciErr = clRetainContext(slot.context);
+       assert(ciErr == CL_SUCCESS);
+       (void)ciErr;
+
+       return slot.context;
+}
+
+cl_program OpenCLCache::get_program(cl_platform_id platform,
+                                    cl_device_id device,
+                                    ustring key,
+                                    thread_scoped_lock& slot_locker)
+{
+       assert(platform != NULL);
+
+       OpenCLCache& self = global_instance();
+
+       thread_scoped_lock cache_lock(self.cache_lock);
+
+       pair<CacheMap::iterator,bool> ins = self.cache.insert(
+               CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
+
+       Slot &slot = ins.first->second;
+
+       pair<Slot::EntryMap::iterator,bool> ins2 = slot.programs.insert(
+               Slot::EntryMap::value_type(key, Slot::ProgramEntry()));
+
+       Slot::ProgramEntry &entry = ins2.first->second;
+
+       /* create slot lock only while holding cache lock */
+       if(!entry.mutex)
+               entry.mutex = new thread_mutex;
+
+       /* need to unlock cache before locking slot, to allow store to complete */
+       cache_lock.unlock();
+
+       /* lock the slot */
+       slot_locker = thread_scoped_lock(*entry.mutex);
+
+       /* If the thing isn't cached */
+       if(entry.program == NULL) {
+               /* return with the caller's lock holder holding the slot lock */
+               return NULL;
+       }
+
+       /* the item was already cached, release the slot lock */
+       slot_locker.unlock();
+
+       cl_int ciErr = clRetainProgram(entry.program);
+       assert(ciErr == CL_SUCCESS);
+       (void)ciErr;
+
+       return entry.program;
+}
+
+void OpenCLCache::store_context(cl_platform_id platform,
+                                cl_device_id device,
+                                cl_context context,
+                                thread_scoped_lock& slot_locker)
+{
+       assert(platform != NULL);
+       assert(device != NULL);
+       assert(context != NULL);
+
+       OpenCLCache &self = global_instance();
+
+       thread_scoped_lock cache_lock(self.cache_lock);
+       CacheMap::iterator i = self.cache.find(PlatformDevicePair(platform, device));
+       cache_lock.unlock();
+
+       Slot &slot = i->second;
+
+       /* sanity check */
+       assert(i != self.cache.end());
+       assert(slot.context == NULL);
+
+       slot.context = context;
+
+       /* unlock the slot */
+       slot_locker.unlock();
+
+       /* increment reference count in OpenCL.
+        * The caller is going to release the object when done with it. */
+       cl_int ciErr = clRetainContext(context);
+       assert(ciErr == CL_SUCCESS);
+       (void)ciErr;
+}
+
+void OpenCLCache::store_program(cl_platform_id platform,
+                                cl_device_id device,
+                                cl_program program,
+                                ustring key,
+                                thread_scoped_lock& slot_locker)
+{
+       assert(platform != NULL);
+       assert(device != NULL);
+       assert(program != NULL);
+
+       OpenCLCache &self = global_instance();
+
+       thread_scoped_lock cache_lock(self.cache_lock);
+
+       CacheMap::iterator i = self.cache.find(PlatformDevicePair(platform, device));
+       assert(i != self.cache.end());
+       Slot &slot = i->second;
+
+       Slot::EntryMap::iterator i2 = slot.programs.find(key);
+       assert(i2 != slot.programs.end());
+       Slot::ProgramEntry &entry = i2->second;
+
+       assert(entry.program == NULL);
+
+       cache_lock.unlock();
+
+       entry.program = program;
+
+       /* unlock the slot */
+       slot_locker.unlock();
+
+       /* Increment reference count in OpenCL.
+        * The caller is going to release the object when done with it.
+        */
+       cl_int ciErr = clRetainProgram(program);
+       assert(ciErr == CL_SUCCESS);
+       (void)ciErr;
+}
+
+string OpenCLCache::get_kernel_md5()
+{
+       OpenCLCache &self = global_instance();
+       thread_scoped_lock lock(self.kernel_md5_lock);
+
+       if(self.kernel_md5.empty()) {
+               self.kernel_md5 = path_files_md5_hash(path_get("kernel"));
+       }
+       return self.kernel_md5;
+}
+
+OpenCLDeviceBase::OpenCLProgram::OpenCLProgram(OpenCLDeviceBase *device, string program_name, string kernel_file, string kernel_build_options)
+ : device(device),
+   program_name(program_name),
+   kernel_file(kernel_file),
+   kernel_build_options(kernel_build_options)
+{
+       loaded = false;
+       program = NULL;
+}
+
+OpenCLDeviceBase::OpenCLProgram::~OpenCLProgram()
+{
+       release();
+}
+
+void OpenCLDeviceBase::OpenCLProgram::release()
+{
+       for(map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end(); ++kernel) {
+               if(kernel->second) {
+                       clReleaseKernel(kernel->second);
+                       kernel->second = NULL;
+               }
+       }
+       if(program) {
+               clReleaseProgram(program);
+               program = NULL;
+       }
+}
+
+void OpenCLDeviceBase::OpenCLProgram::add_kernel(ustring name)
+{
+       if(!kernels.count(name)) {
+               kernels[name] = NULL;
+       }
+}
+
+bool OpenCLDeviceBase::OpenCLProgram::build_kernel(const string *debug_src)
+{
+       string build_options;
+       build_options = device->kernel_build_options(debug_src) + kernel_build_options;
+
+       cl_int ciErr = clBuildProgram(program, 0, NULL, build_options.c_str(), NULL, NULL);
+
+       /* show warnings even if build is successful */
+       size_t ret_val_size = 0;
+
+       clGetProgramBuildInfo(program, device->cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+
+       if(ret_val_size > 1) {
+               vector<char> build_log(ret_val_size + 1);
+               clGetProgramBuildInfo(program, device->cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, &build_log[0], NULL);
+
+               build_log[ret_val_size] = '\0';
+               /* Skip meaningless empty output from the NVidia compiler. */
+               if(!(ret_val_size == 2 && build_log[0] == '\n')) {
+                       output_msg = string(&build_log[0]);
+               }
+       }
+
+       if(ciErr != CL_SUCCESS) {
+               error_msg = string("OpenCL build failed: ") + clewErrorString(ciErr);
+               return false;
+       }
+
+       return true;
+}
+
+bool OpenCLDeviceBase::OpenCLProgram::compile_kernel(const string *debug_src)
+{
+       string source = "#include \"kernels/opencl/" + kernel_file + "\" // " + OpenCLCache::get_kernel_md5() + "\n";
+       /* We compile kernels consisting of many files. unfortunately OpenCL
+        * kernel caches do not seem to recognize changes in included files.
+        * so we force recompile on changes by adding the md5 hash of all files.
+        */
+       source = path_source_replace_includes(source, path_get("kernel"));
+
+       if(debug_src) {
+               path_write_text(*debug_src, source);
+       }
+
+       size_t source_len = source.size();
+       const char *source_str = source.c_str();
+       cl_int ciErr;
+
+       program = clCreateProgramWithSource(device->cxContext,
+                                          1,
+                                          &source_str,
+                                          &source_len,
+                                          &ciErr);
+
+       if(ciErr != CL_SUCCESS) {
+               error_msg = string("OpenCL program creation failed: ") + clewErrorString(ciErr);
+               return false;
+       }
+
+       double starttime = time_dt();
+
+       log += "Build flags: " + kernel_build_options + "\n";
+
+       if(!build_kernel(debug_src))
+               return false;
+
+       log += "Kernel compilation of " + program_name + " finished in " + string_printf("%.2lfs.\n", time_dt() - starttime);
+
+       return true;
+}
+
+bool OpenCLDeviceBase::OpenCLProgram::load_binary(const string& clbin,
+                                                  const string *debug_src)
+{
+       /* read binary into memory */
+       vector<uint8_t> binary;
+
+       if(!path_read_binary(clbin, binary)) {
+               error_msg = "OpenCL failed to read cached binary " + clbin + ".";
+               return false;
+       }
+
+       /* create program */
+       cl_int status, ciErr;
+       size_t size = binary.size();
+       const uint8_t *bytes = &binary[0];
+
+       program = clCreateProgramWithBinary(device->cxContext, 1, &device->cdDevice,
+               &size, &bytes, &status, &ciErr);
+
+       if(status != CL_SUCCESS || ciErr != CL_SUCCESS) {
+               error_msg = "OpenCL failed create program from cached binary " + clbin + ": " + clewErrorString(status) + " " + clewErrorString(ciErr);
+               return false;
+       }
+
+       if(!build_kernel(debug_src))
+               return false;
+
+       return true;
+}
+
+bool OpenCLDeviceBase::OpenCLProgram::save_binary(const string& clbin)
+{
+       size_t size = 0;
+       clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
+
+       if(!size)
+               return false;
+
+       vector<uint8_t> binary(size);
+       uint8_t *bytes = &binary[0];
+
+       clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
+
+       return path_write_binary(clbin, binary);
+}
+
+void OpenCLDeviceBase::OpenCLProgram::load()
+{
+       assert(device);
+
+       loaded = false;
+
+       string device_md5 = device->device_md5_hash(kernel_build_options);
+
+       /* Try to use cached kernel. */
+       thread_scoped_lock cache_locker;
+       ustring cache_key(program_name + device_md5);
+       program = device->load_cached_kernel(cache_key,
+                                            cache_locker);
+
+       if(!program) {
+               log += "OpenCL program " + program_name + " not found in cache.\n";
+
+               string basename = "cycles_kernel_" + program_name + "_" + device_md5 + "_" + OpenCLCache::get_kernel_md5();
+               basename = path_cache_get(path_join("kernels", basename));
+               string clbin = basename + ".clbin";
+
+