diff options
Diffstat (limited to 'intern/cycles/device/device_opencl.cpp')
-rw-r--r-- | intern/cycles/device/device_opencl.cpp | 3268 |
1 files changed, 11 insertions, 3257 deletions
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index 830e4d056b5..45cf6b074e9 100644 --- a/intern/cycles/device/device_opencl.cpp +++ b/intern/cycles/device/device_opencl.cpp @@ -16,3275 +16,29 @@ #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 = ""; |